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

aelovikov-intel · web-flow · commit 4ab1e4b9eccd · 2024-07-08T11:08:32.000+02:00
Not needed anymore after deprecated shuffles/collectives were removed.
diff --git a/sycl/include/sycl/detail/group_sort_impl.hpp b/sycl/include/sycl/detail/group_sort_impl.hpp
@@ -12,6 +12,8 @@
 
 #ifdef __SYCL_DEVICE_ONLY__
 
+#include <climits>
+
 #include <sycl/builtins.hpp>
 #include <sycl/group_algorithm.hpp>
 #include <sycl/group_barrier.hpp>
diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp
@@ -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>
diff --git a/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp
@@ -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
diff --git a/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp
@@ -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>
diff --git a/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp
@@ -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>
@@ -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
 
diff --git a/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
@@ -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 {
diff --git a/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp
@@ -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>
diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
@@ -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>
diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp
@@ -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
 
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
@@ -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
@@ -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__
   }
@@ -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__
   }
diff --git a/sycl/test-e2e/Assert/check_resource_leak.cpp b/sycl/test-e2e/Assert/check_resource_leak.cpp
@@ -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>
 
diff --git a/sycl/test-e2e/Basic/image/image.cpp b/sycl/test-e2e/Basic/image/image.cpp
@@ -12,6 +12,7 @@
 //===----------------------------------------------------------------------===//
 
 #include <sycl/accessor_image.hpp>
+#include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
 
 #include <iostream>
diff --git a/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp b/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp
@@ -15,6 +15,7 @@
 //===----------------------------------------------------------------------===//
 
 #include <sycl/accessor_image.hpp>
+#include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
 #include <sycl/image.hpp>
 
diff --git a/sycl/test-e2e/Basic/image/image_array.cpp b/sycl/test-e2e/Basic/image/image_array.cpp
@@ -15,6 +15,7 @@
 //===----------------------------------------------------------------------===//
 
 #include <sycl/accessor_image.hpp>
+#include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
 
 #include <iostream>
diff --git a/sycl/test-e2e/DeviceLib/assert.cpp b/sycl/test-e2e/DeviceLib/assert.cpp
@@ -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>
diff --git a/sycl/test-e2e/DeviceLib/rand_test.cpp b/sycl/test-e2e/DeviceLib/rand_test.cpp
@@ -3,6 +3,7 @@
 
 // UNSUPPORTED: cuda || hip
 
+#include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
 
 #include <cstdlib>
diff --git a/sycl/test/check_device_code/extensions/sub_group_as.cpp b/sycl/test/check_device_code/extensions/sub_group_as.cpp
@@ -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>
@@ -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)
diff --git a/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp b/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp
@@ -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;
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
@@ -133,31 +133,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: event.hpp
@@ -183,4 +158,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:

Original file line number	Diff line number	Diff line change
`@@ -52,6 +52,9 @@ GetMultiPtrDecoratedAs(multi_ptr<FromT, Space, IsDecorated> MPtr) {`
`52`	`52`	`MPtr.get_decorated());`
`53`	`53`	`}`
`54`	`54`
	`55`	`+template <typename NonUniformGroup>`
	`56`	`+inline uint32_t IdToMaskPosition(NonUniformGroup Group, uint32_t Id);`
	`57`	`+`
`55`	`58`	`namespace spirv {`
`56`	`59`
`57`	`60`	`template <typename Group>`