Skip to content

[SYCL] Drop sycl/ext/oneapi/functional.hpp include from sycl/sub_group.hpp #13760

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jul 8, 2024
Merged

[SYCL] Drop sycl/ext/oneapi/functional.hpp include from sycl/sub_group.hpp #13760

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions sycl/include/sycl/detail/group_sort_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,8 @@

#ifdef __SYCL_DEVICE_ONLY__

#include <climits>

#include <sycl/builtins.hpp>
#include <sycl/group_algorithm.hpp>
#include <sycl/group_barrier.hpp>
Expand Down
3 changes: 3 additions & 0 deletions sycl/include/sycl/detail/spirv.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,6 +52,9 @@ GetMultiPtrDecoratedAs(multi_ptr<FromT, Space, IsDecorated> MPtr) {
MPtr.get_decorated());
}

template <typename NonUniformGroup>
inline uint32_t IdToMaskPosition(NonUniformGroup Group, uint32_t Id);

namespace spirv {

template <typename Group>
Expand Down
3 changes: 2 additions & 1 deletion sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,8 @@
#pragma once

#include <sycl/aspects.hpp>
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
#include <sycl/detail/spirv.hpp>
#include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
#include <sycl/exception.hpp> // for runtime_error
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp> // for GetMask
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@

#include <sycl/aspects.hpp>
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
#include <sycl/detail/spirv.hpp>
#include <sycl/detail/type_traits.hpp> // for is_fixed_size_group, is_group
#include <sycl/exception.hpp> // for runtime_error
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
Expand Down
5 changes: 5 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@

#include <sycl/aspects.hpp>
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
#include <sycl/detail/spirv.hpp>
#include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
#include <sycl/exception.hpp> // for runtime_error
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
Expand All @@ -20,6 +21,10 @@
#include <sycl/range.hpp> // for range
#include <sycl/sub_group.hpp>

#ifdef __SYCL_DEVICE_ONLY__
#include <sycl/ext/oneapi/functional.hpp>
#endif

#include <stdint.h> // for uint32_t
#include <type_traits> // for true_type

Expand Down
5 changes: 5 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,13 +8,18 @@

#pragma once

#include <sycl/detail/spirv.hpp>
#include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp>
#include <sycl/ext/oneapi/free_function_queries.hpp>
#include <sycl/group.hpp>
#include <sycl/memory_enums.hpp>
#include <sycl/nd_item.hpp>
#include <sycl/sub_group.hpp>

#ifdef __SYCL_DEVICE_ONLY__
#include <sycl/ext/oneapi/functional.hpp>
#endif

namespace sycl {
inline namespace _V1 {
namespace ext::oneapi::experimental {
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@

#include <sycl/aspects.hpp>
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
#include <sycl/detail/spirv.hpp>
#include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
#include <sycl/exception.hpp> // for runtime_error
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
Expand Down
8 changes: 8 additions & 0 deletions sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -359,3 +359,11 @@ group_ballot(Group g, bool predicate) {
} // namespace ext::oneapi
} // namespace _V1
} // namespace sycl

// We have a cyclic dependency with
// sub_group_mask.hpp
// detail/spirv.hpp
// non_uniform_groups.hpp
// "Break" it by including this at the end (instead of beginning). Ideally, we
// should refactor this somehow...
#include <sycl/detail/spirv.hpp>
1 change: 1 addition & 0 deletions sycl/include/sycl/group_algorithm.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@
#include <sycl/ext/oneapi/functional.hpp>
#if defined(__NVPTX__)
#include <sycl/ext/oneapi/experimental/cuda/non_uniform_algorithms.hpp>
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
#endif
#endif

Expand Down
9 changes: 2 additions & 7 deletions sycl/include/sycl/sub_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,10 +19,6 @@
#include <sycl/multi_ptr.hpp> // for multi_ptr
#include <sycl/range.hpp> // for range

#ifdef __SYCL_DEVICE_ONLY__
#include <sycl/ext/oneapi/functional.hpp>
#endif

#include <stdint.h> // for uint32_t
#include <tuple> // for _Swallow_assign, ignore
#include <type_traits> // for enable_if_t, remove_cv_t
Expand Down Expand Up @@ -239,7 +235,7 @@ struct sub_group {
if (g)
return load(g);

assert(!"Sub-group load() is supported for local or global pointers only.");
// Sub-group load() is supported for local or global pointers only.
return {};
#endif // __NVPTX__ || __AMDGCN__
}
Expand Down Expand Up @@ -421,8 +417,7 @@ struct sub_group {
return;
}

assert(
!"Sub-group store() is supported for local or global pointers only.");
// Sub-group store() is supported for local or global pointers only.
return;
#endif // __NVPTX__ || __AMDGCN__
}
Expand Down
2 changes: 2 additions & 0 deletions sycl/test-e2e/Assert/check_resource_leak.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,8 @@
#include <sycl/detail/core.hpp>

#include <sycl/builtins.hpp>
#include <sycl/detail/spirv.hpp>
#include <sycl/ext/oneapi/experimental/ballot_group.hpp>
#include <sycl/ext/oneapi/free_function_queries.hpp>
#include <sycl/ext/oneapi/sub_group_mask.hpp>

Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/Basic/image/image.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,7 @@
//===----------------------------------------------------------------------===//

#include <sycl/accessor_image.hpp>
#include <sycl/builtins.hpp>
#include <sycl/detail/core.hpp>

#include <iostream>
Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@
//===----------------------------------------------------------------------===//

#include <sycl/accessor_image.hpp>
#include <sycl/builtins.hpp>
#include <sycl/detail/core.hpp>
#include <sycl/image.hpp>

Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/Basic/image/image_array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@
//===----------------------------------------------------------------------===//

#include <sycl/accessor_image.hpp>
#include <sycl/builtins.hpp>
#include <sycl/detail/core.hpp>

#include <iostream>
Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/DeviceLib/assert.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,6 +95,7 @@
#include <assert.h>
#include <iostream>
#include <stdlib.h>
#include <sycl/builtins.hpp>
#include <sycl/detail/core.hpp>
#include <sys/types.h>
#include <sys/wait.h>
Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/DeviceLib/rand_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,7 @@

// UNSUPPORTED: cuda || hip

#include <sycl/builtins.hpp>
#include <sycl/detail/core.hpp>

#include <cstdlib>
Expand Down
7 changes: 0 additions & 7 deletions sycl/test/check_device_code/extensions/sub_group_as.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,6 @@
// RUN: %clangxx -fsycl-device-only -O0 -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O0
// Test compilation with -O3 when all methods are inlined in kernel function
// and -O0 when helper methods are preserved.
#include <cassert>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
Expand Down Expand Up @@ -46,45 +45,39 @@ SYCL_EXTERNAL void test(sycl::accessor<int, 1, sycl::access::mode::read_write,
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
// CHECK-O3: call spir_func void {{.*}}assert


// load() for local address space
// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
// CHECK-O3: call spir_func void {{.*}}assert

// load() for private address space
// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
// CHECK-O3: call spir_func void {{.*}}assert

// store() for global address space
// NOTE: Call to __spirv_GenericCastToPtrExplicit_ToLocal is consolidated with an earlier call to it.
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
// CHECK-O3: call spir_func void {{.*}}spirv_SubgroupBlockWriteINTEL{{.*}}(ptr addrspace(1)
// CHECK-O3: call spir_func void {{.*}}assert

// load() accepting raw pointers method
// CHECK-O0: define{{.*}}spir_func i32 {{.*}}4sycl3_V19sub_group4load{{.*}}addrspace(4) %
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
// CHECK-O0: call spir_func i32 {{.*}}sycl3_V19sub_group4load{{.*}}ptr addrspace(3) %
// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
// CHECK-O0: call spir_func i32 {{.*}}sycl3_V19sub_group4load{{.*}}ptr addrspace(1) %
// CHECK-O0: call spir_func void {{.*}}assert

// store() accepting raw pointers method
// CHECK-O0: define{{.*}}spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}ptr addrspace(4) %
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
// CHECK-O0: call spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}, ptr addrspace(3) %
// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
// CHECK-O0: call spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}, ptr addrspace(1) %
// CHECK-O0: call spir_func void {{.*}}assert

// CHECK-O0: define {{.*}}spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4) %
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
Expand Down
1 change: 1 addition & 0 deletions sycl/test/check_device_code/vector/vector_convert_bfloat.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@

#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bfloat16.hpp>
#include <sycl/types.hpp>

using namespace sycl;
using bfloat16 = sycl::ext::oneapi::bfloat16;
Expand Down
26 changes: 1 addition & 25 deletions sycl/test/include_deps/sycl_detail_core.hpp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -134,31 +134,6 @@
// CHECK-NEXT: nd_item.hpp
// CHECK-NEXT: nd_range.hpp
// CHECK-NEXT: sub_group.hpp
// CHECK-NEXT: ext/oneapi/functional.hpp
// CHECK-NEXT: detail/spirv.hpp
// CHECK-NEXT: ext/oneapi/experimental/non_uniform_groups.hpp
// CHECK-NEXT: ext/oneapi/sub_group_mask.hpp
// CHECK-NEXT: builtins.hpp
// CHECK-NEXT: detail/builtins/builtins.hpp
// CHECK-NEXT: builtins_utils_vec.hpp
// CHECK-NEXT: builtins_utils_scalar.hpp
// CHECK-NEXT: detail/boolean.hpp
// CHECK-NEXT: marray.hpp
// CHECK-NEXT: types.hpp
// CHECK-NEXT: vector.hpp
// CHECK-NEXT: detail/memcpy.hpp
// CHECK-NEXT: detail/vector_convert.hpp
// CHECK-NEXT: swizzles.def
// CHECK-NEXT: detail/builtins/common_functions.inc
// CHECK-NEXT: detail/builtins/helper_macros.hpp
// CHECK-NEXT: detail/builtins/geometric_functions.inc
// CHECK-NEXT: detail/builtins/half_precision_math_functions.inc
// CHECK-NEXT: detail/builtins/integer_functions.inc
// CHECK-NEXT: detail/builtins/math_functions.inc
// CHECK-NEXT: detail/builtins/native_math_functions.inc
// CHECK-NEXT: detail/builtins/relational_functions.inc
// CHECK-NEXT: feature_test.hpp
// CHECK-NEXT: functional.hpp
// CHECK-NEXT: device.hpp
// CHECK-NEXT: kernel_bundle_enums.hpp
// CHECK-NEXT: exception_list.hpp
Expand All @@ -181,4 +156,5 @@
// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
// CHECK-NEXT: feature_test.hpp
// CHECK-EMPTY:
Loading