Skip to content

Commit 903ac62

Browse files
aelovikov-intelaelovikov-intel
committed
[SYCL] Drop sycl/ext/oneapi/functional.hpp include from sycl/sub_group.hpp
Not needed anymore after deprecated shuffles/collectives were removed. I had to remove usage of `assert` from `sub_group`'s `load`/`store` methods to avoid the last dependency on `sycl/builtins.hpp` from `sycl/detail/core.hpp`. That should be fine for two reasons: 1) These are extension methods and not part of the SYCL 2020 2) `assert` support isn't universal on devices anyway, and fallback support is expensive and is disabled by default, so we are not using them in device code generally.
1 parent 4a87b2c commit 903ac62

19 files changed

+39
-40
lines changed

‎sycl/include/sycl/detail/group_sort_impl.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,8 @@
1212

1313
#ifdef __SYCL_DEVICE_ONLY__
1414

15+
#include <climits>
16+
1517
#include <sycl/builtins.hpp>
1618
#include <sycl/group_algorithm.hpp>
1719
#include <sycl/group_barrier.hpp>

‎sycl/include/sycl/detail/spirv.hpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,9 @@ GetMultiPtrDecoratedAs(multi_ptr<FromT, Space, IsDecorated> MPtr) {
5252
MPtr.get_decorated());
5353
}
5454

55+
template <typename NonUniformGroup>
56+
inline uint32_t IdToMaskPosition(NonUniformGroup Group, uint32_t Id);
57+
5558
namespace spirv {
5659

5760
template <typename Group>

‎sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,8 @@
99
#pragma once
1010

1111
#include <sycl/aspects.hpp>
12-
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
12+
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
13+
#include <sycl/detail/spirv.hpp>
1314
#include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
1415
#include <sycl/exception.hpp> // for runtime_error
1516
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp> // for GetMask

‎sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010

1111
#include <sycl/aspects.hpp>
1212
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
13+
#include <sycl/detail/spirv.hpp>
1314
#include <sycl/detail/type_traits.hpp> // for is_fixed_size_group, is_group
1415
#include <sycl/exception.hpp> // for runtime_error
1516
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>

‎sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010

1111
#include <sycl/aspects.hpp>
1212
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
13+
#include <sycl/detail/spirv.hpp>
1314
#include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
1415
#include <sycl/exception.hpp> // for runtime_error
1516
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
@@ -20,6 +21,10 @@
2021
#include <sycl/range.hpp> // for range
2122
#include <sycl/sub_group.hpp>
2223

24+
#ifdef __SYCL_DEVICE_ONLY__
25+
#include <sycl/ext/oneapi/functional.hpp>
26+
#endif
27+
2328
#include <stdint.h> // for uint32_t
2429
#include <type_traits> // for true_type
2530

‎sycl/include/sycl/ext/oneapi/experimental/root_group.hpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,13 +8,18 @@
88

99
#pragma once
1010

11+
#include <sycl/detail/spirv.hpp>
1112
#include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp>
1213
#include <sycl/ext/oneapi/free_function_queries.hpp>
1314
#include <sycl/group.hpp>
1415
#include <sycl/memory_enums.hpp>
1516
#include <sycl/nd_item.hpp>
1617
#include <sycl/sub_group.hpp>
1718

19+
#ifdef __SYCL_DEVICE_ONLY__
20+
#include <sycl/ext/oneapi/functional.hpp>
21+
#endif
22+
1823
namespace sycl {
1924
inline namespace _V1 {
2025
namespace ext::oneapi::experimental {

‎sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010

1111
#include <sycl/aspects.hpp>
1212
#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
13+
#include <sycl/detail/spirv.hpp>
1314
#include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
1415
#include <sycl/exception.hpp> // for runtime_error
1516
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>

‎sycl/include/sycl/ext/oneapi/sub_group_mask.hpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -359,3 +359,11 @@ group_ballot(Group g, bool predicate) {
359359
} // namespace ext::oneapi
360360
} // namespace _V1
361361
} // namespace sycl
362+
363+
// We have a cyclic dependency with
364+
// sub_group_mask.hpp
365+
// detail/spirv.hpp
366+
// non_uniform_groups.hpp
367+
// "Break" it by including this at the end (instead of beginning). Ideally, we
368+
// should refactor this somehow...
369+
#include <sycl/detail/spirv.hpp>

‎sycl/include/sycl/group_algorithm.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -28,6 +28,7 @@
2828
#include <sycl/ext/oneapi/functional.hpp>
2929
#if defined(__NVPTX__)
3030
#include <sycl/ext/oneapi/experimental/cuda/non_uniform_algorithms.hpp>
31+
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
3132
#endif
3233
#endif
3334

‎sycl/include/sycl/sub_group.hpp

Lines changed: 2 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -19,10 +19,6 @@
1919
#include <sycl/multi_ptr.hpp> // for multi_ptr
2020
#include <sycl/range.hpp> // for range
2121

22-
#ifdef __SYCL_DEVICE_ONLY__
23-
#include <sycl/ext/oneapi/functional.hpp>
24-
#endif
25-
2622
#include <stdint.h> // for uint32_t
2723
#include <tuple> // for _Swallow_assign, ignore
2824
#include <type_traits> // for enable_if_t, remove_cv_t
@@ -239,7 +235,7 @@ struct sub_group {
239235
if (g)
240236
return load(g);
241237

242-
assert(!"Sub-group load() is supported for local or global pointers only.");
238+
// Sub-group load() is supported for local or global pointers only.
243239
return {};
244240
#endif // __NVPTX__ || __AMDGCN__
245241
}
@@ -421,8 +417,7 @@ struct sub_group {
421417
return;
422418
}
423419

424-
assert(
425-
!"Sub-group store() is supported for local or global pointers only.");
420+
// Sub-group store() is supported for local or global pointers only.
426421
return;
427422
#endif // __NVPTX__ || __AMDGCN__
428423
}

‎sycl/test-e2e/Assert/check_resource_leak.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,8 @@
1111
#include <sycl/detail/core.hpp>
1212

1313
#include <sycl/builtins.hpp>
14+
#include <sycl/detail/spirv.hpp>
15+
#include <sycl/ext/oneapi/experimental/ballot_group.hpp>
1416
#include <sycl/ext/oneapi/free_function_queries.hpp>
1517
#include <sycl/ext/oneapi/sub_group_mask.hpp>
1618

‎sycl/test-e2e/Basic/image/image.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,7 @@
1212
//===----------------------------------------------------------------------===//
1313

1414
#include <sycl/accessor_image.hpp>
15+
#include <sycl/builtins.hpp>
1516
#include <sycl/detail/core.hpp>
1617

1718
#include <iostream>

‎sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@
1515
//===----------------------------------------------------------------------===//
1616

1717
#include <sycl/accessor_image.hpp>
18+
#include <sycl/builtins.hpp>
1819
#include <sycl/detail/core.hpp>
1920
#include <sycl/image.hpp>
2021

‎sycl/test-e2e/Basic/image/image_array.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@
1515
//===----------------------------------------------------------------------===//
1616

1717
#include <sycl/accessor_image.hpp>
18+
#include <sycl/builtins.hpp>
1819
#include <sycl/detail/core.hpp>
1920

2021
#include <iostream>

‎sycl/test-e2e/DeviceLib/assert.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -95,6 +95,7 @@
9595
#include <assert.h>
9696
#include <iostream>
9797
#include <stdlib.h>
98+
#include <sycl/builtins.hpp>
9899
#include <sycl/detail/core.hpp>
99100
#include <sys/types.h>
100101
#include <sys/wait.h>

‎sycl/test-e2e/DeviceLib/rand_test.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,7 @@
33

44
// UNSUPPORTED: cuda || hip
55

6+
#include <sycl/builtins.hpp>
67
#include <sycl/detail/core.hpp>
78

89
#include <cstdlib>

‎sycl/test/check_device_code/extensions/sub_group_as.cpp

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,6 @@
22
// RUN: %clangxx -fsycl-device-only -O0 -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O0
33
// Test compilation with -O3 when all methods are inlined in kernel function
44
// and -O0 when helper methods are preserved.
5-
#include <cassert>
65
#include <cstdint>
76
#include <cstdio>
87
#include <cstdlib>
@@ -46,45 +45,39 @@ SYCL_EXTERNAL void test(sycl::accessor<int, 1, sycl::access::mode::read_write,
4645
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
4746
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
4847
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
49-
// CHECK-O3: call spir_func void {{.*}}assert
5048

5149

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

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

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

7368
// load() accepting raw pointers method
7469
// CHECK-O0: define{{.*}}spir_func i32 {{.*}}4sycl3_V19sub_group4load{{.*}}addrspace(4) %
7570
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
7671
// CHECK-O0: call spir_func i32 {{.*}}sycl3_V19sub_group4load{{.*}}ptr addrspace(3) %
7772
// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
7873
// CHECK-O0: call spir_func i32 {{.*}}sycl3_V19sub_group4load{{.*}}ptr addrspace(1) %
79-
// CHECK-O0: call spir_func void {{.*}}assert
8074

8175
// store() accepting raw pointers method
8276
// CHECK-O0: define{{.*}}spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}ptr addrspace(4) %
8377
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
8478
// CHECK-O0: call spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}, ptr addrspace(3) %
8579
// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
8680
// CHECK-O0: call spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}, ptr addrspace(1) %
87-
// CHECK-O0: call spir_func void {{.*}}assert
8881

8982
// CHECK-O0: define {{.*}}spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4) %
9083
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)

‎sycl/test/check_device_code/vector/vector_convert_bfloat.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66

77
#include <sycl/detail/core.hpp>
88
#include <sycl/ext/oneapi/bfloat16.hpp>
9+
#include <sycl/types.hpp>
910

1011
using namespace sycl;
1112
using bfloat16 = sycl::ext::oneapi::bfloat16;

‎sycl/test/include_deps/sycl_detail_core.hpp.cpp

Lines changed: 1 addition & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -134,31 +134,6 @@
134134
// CHECK-NEXT: nd_item.hpp
135135
// CHECK-NEXT: nd_range.hpp
136136
// CHECK-NEXT: sub_group.hpp
137-
// CHECK-NEXT: ext/oneapi/functional.hpp
138-
// CHECK-NEXT: detail/spirv.hpp
139-
// CHECK-NEXT: ext/oneapi/experimental/non_uniform_groups.hpp
140-
// CHECK-NEXT: ext/oneapi/sub_group_mask.hpp
141-
// CHECK-NEXT: builtins.hpp
142-
// CHECK-NEXT: detail/builtins/builtins.hpp
143-
// CHECK-NEXT: builtins_utils_vec.hpp
144-
// CHECK-NEXT: builtins_utils_scalar.hpp
145-
// CHECK-NEXT: detail/boolean.hpp
146-
// CHECK-NEXT: marray.hpp
147-
// CHECK-NEXT: types.hpp
148-
// CHECK-NEXT: vector.hpp
149-
// CHECK-NEXT: detail/memcpy.hpp
150-
// CHECK-NEXT: detail/vector_convert.hpp
151-
// CHECK-NEXT: swizzles.def
152-
// CHECK-NEXT: detail/builtins/common_functions.inc
153-
// CHECK-NEXT: detail/builtins/helper_macros.hpp
154-
// CHECK-NEXT: detail/builtins/geometric_functions.inc
155-
// CHECK-NEXT: detail/builtins/half_precision_math_functions.inc
156-
// CHECK-NEXT: detail/builtins/integer_functions.inc
157-
// CHECK-NEXT: detail/builtins/math_functions.inc
158-
// CHECK-NEXT: detail/builtins/native_math_functions.inc
159-
// CHECK-NEXT: detail/builtins/relational_functions.inc
160-
// CHECK-NEXT: feature_test.hpp
161-
// CHECK-NEXT: functional.hpp
162137
// CHECK-NEXT: device.hpp
163138
// CHECK-NEXT: kernel_bundle_enums.hpp
164139
// CHECK-NEXT: exception_list.hpp
@@ -181,4 +156,5 @@
181156
// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
182157
// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
183158
// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
159+
// CHECK-NEXT: feature_test.hpp
184160
// CHECK-EMPTY:

0 commit comments

Comments
 (0)