[SYCL] Fix post-commit issues from ballot_group (#9164)

- Remove __builtin_unreachable()
- Suppress warnings about unused parameters for host device
- Rename cluster_group in end-to-end tests (see #8978)

---------

Signed-off-by: John Pennycook <[email protected]>
Co-authored-by: Larsen, Steffen <[email protected]>

Author: Pennycook

sycl/include/sycl/ext/oneapi/experimental/non_uniform_groups.hpp

@@ -54,7 +54,6 @@ inline uint32_t IdToMaskPosition(NonUniformGroup Group, uint32_t Id) {
       }
     }
   }
-  __builtin_unreachable();
   return Count;
 }
 

sycl/include/sycl/ext/oneapi/group_algorithm.hpp

@@ -151,6 +151,7 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::spirv::GroupBroadcast(g, x, local_id);
 #else
+  (void)g;
   (void)x;
   (void)local_id;
   throw runtime_error("Group algorithms are not supported on host device.",

sycl/include/sycl/group_algorithm.hpp

@@ -208,6 +208,7 @@ reduce_over_group(Group g, T x, BinaryOperation binary_op) {
   return sycl::detail::calc<__spv::GroupOperation::Reduce>(
       g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
 #else
+  (void)g;
   throw runtime_error("Group algorithms are not supported on host.",
                       PI_ERROR_INVALID_DEVICE);
 #endif
@@ -379,6 +380,7 @@ any_of_group(Group g, bool pred) {
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::spirv::GroupAny(g, pred);
 #else
+  (void)g;
   (void)pred;
   throw runtime_error("Group algorithms are not supported on host.",
                       PI_ERROR_INVALID_DEVICE);
@@ -418,6 +420,7 @@ all_of_group(Group g, bool pred) {
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::spirv::GroupAll(g, pred);
 #else
+  (void)g;
   (void)pred;
   throw runtime_error("Group algorithms are not supported on host.",
                       PI_ERROR_INVALID_DEVICE);
@@ -457,6 +460,7 @@ none_of_group(Group g, bool pred) {
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::spirv::GroupAll(g, !pred);
 #else
+  (void)g;
   (void)pred;
   throw runtime_error("Group algorithms are not supported on host.",
                       PI_ERROR_INVALID_DEVICE);
@@ -574,6 +578,7 @@ group_broadcast(Group g, T x, typename Group::id_type local_id) {
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::spirv::GroupBroadcast(g, x, local_id);
 #else
+  (void)g;
   (void)x;
   (void)local_id;
   throw runtime_error("Group algorithms are not supported on host.",
@@ -637,6 +642,7 @@ exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
   return sycl::detail::calc<__spv::GroupOperation::ExclusiveScan>(
       g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
 #else
+  (void)g;
   throw runtime_error("Group algorithms are not supported on host.",
                       PI_ERROR_INVALID_DEVICE);
 #endif
@@ -870,6 +876,7 @@ inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
   return sycl::detail::calc<__spv::GroupOperation::InclusiveScan>(
       g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
 #else
+  (void)g;
   throw runtime_error("Group algorithms are not supported on host.",
                       PI_ERROR_INVALID_DEVICE);
 #endif

sycl/include/sycl/group_barrier.hpp

@@ -29,6 +29,8 @@ group_barrier(Group G, memory_scope FenceScope = Group::fence_scope) {
 #ifdef __SYCL_DEVICE_ONLY__
   detail::spirv::ControlBarrier(G, FenceScope, memory_order::seq_cst);
 #else
+  (void)G;
+  (void)FenceScope;
   throw sycl::runtime_error("Barriers are not supported on host device",
                             PI_ERROR_INVALID_DEVICE);
 #endif

sycl/test-e2e/NonUniformGroups/cluster_group.cpp renamed to sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp

@@ -7,9 +7,9 @@
 #include <vector>
 namespace syclex = sycl::ext::oneapi::experimental;
 
-template <size_t ClusterSize> class TestKernel;
+template <size_t PartitionSize> class TestKernel;
 
-template <size_t ClusterSize> void test() {
+template <size_t PartitionSize> void test() {
   sycl::queue Q;
 
   auto SGSizes = Q.get_device().get_info<sycl::info::device::sub_group_sizes>();
@@ -30,24 +30,24 @@ template <size_t ClusterSize> void test() {
           auto WI = item.get_global_id();
           auto SG = item.get_sub_group();
 
-          auto ClusterGroup = syclex::get_cluster_group<ClusterSize>(SG);
+          auto Partition = syclex::get_fixed_size_group<PartitionSize>(SG);
 
           bool Match = true;
-          Match &= (ClusterGroup.get_group_id() == (WI / ClusterSize));
-          Match &= (ClusterGroup.get_local_id() == (WI % ClusterSize));
-          Match &= (ClusterGroup.get_group_range() == (32 / ClusterSize));
-          Match &= (ClusterGroup.get_local_range() == ClusterSize);
+          Match &= (Partition.get_group_id() == (WI / PartitionSize));
+          Match &= (Partition.get_local_id() == (WI % PartitionSize));
+          Match &= (Partition.get_group_range() == (32 / PartitionSize));
+          Match &= (Partition.get_local_range() == PartitionSize);
           MatchAcc[WI] = Match;
-          LeaderAcc[WI] = ClusterGroup.leader();
+          LeaderAcc[WI] = Partition.leader();
         };
-    CGH.parallel_for<TestKernel<ClusterSize>>(NDR, KernelFunc);
+    CGH.parallel_for<TestKernel<PartitionSize>>(NDR, KernelFunc);
   });
 
   sycl::host_accessor MatchAcc{MatchBuf, sycl::read_only};
   sycl::host_accessor LeaderAcc{LeaderBuf, sycl::read_only};
   for (int WI = 0; WI < 32; ++WI) {
     assert(MatchAcc[WI] == true);
-    assert(LeaderAcc[WI] == ((WI % ClusterSize) == 0));
+    assert(LeaderAcc[WI] == ((WI % PartitionSize) == 0));
   }
 }
 

sycl/test-e2e/NonUniformGroups/is_user_constructed.cpp

@@ -6,9 +6,9 @@ namespace syclex = sycl::ext::oneapi::experimental;
 static_assert(
     syclex::is_user_constructed_group_v<syclex::ballot_group<sycl::sub_group>>);
 static_assert(syclex::is_user_constructed_group_v<
-              syclex::cluster_group<1, sycl::sub_group>>);
+              syclex::fixed_size_group<1, sycl::sub_group>>);
 static_assert(syclex::is_user_constructed_group_v<
-              syclex::cluster_group<2, sycl::sub_group>>);
+              syclex::fixed_size_group<2, sycl::sub_group>>);
 static_assert(
     syclex::is_user_constructed_group_v<syclex::tangle_group<sycl::sub_group>>);
 static_assert(syclex::is_user_constructed_group_v<syclex::opportunistic_group>);