[SYCL] Fixed-size groups and partitions are renamed to "chunks"

sycl/include/sycl/detail/spirv.hpp

@@ -26,7 +26,7 @@ namespace oneapi {
 struct sub_group;
 namespace experimental {
 template <typename ParentGroup> class ballot_group;
-template <size_t PartitionSize, typename ParentGroup> class fixed_size_group;
+template <size_t ChunkSize, typename ParentGroup> class chunk;
 template <int Dimensions> class root_group;
 template <typename ParentGroup> class tangle_group;
 class opportunistic_group;
@@ -78,9 +78,9 @@ struct is_ballot_group<
 
 template <typename Group> struct is_fixed_size_group : std::false_type {};
 
-template <size_t PartitionSize, typename ParentGroup>
-struct is_fixed_size_group<sycl::ext::oneapi::experimental::fixed_size_group<
-    PartitionSize, ParentGroup>> : std::true_type {};
+template <size_t ChunkSize, typename ParentGroup>
+struct is_chunk<sycl::ext::oneapi::experimental::chunk<
+    ChunkSize, ParentGroup>> : std::true_type {};
 
 template <typename Group> struct group_scope {};
 
@@ -105,9 +105,9 @@ struct group_scope<sycl::ext::oneapi::experimental::ballot_group<ParentGroup>> {
   static constexpr __spv::Scope::Flag value = group_scope<ParentGroup>::value;
 };
 
-template <size_t PartitionSize, typename ParentGroup>
-struct group_scope<sycl::ext::oneapi::experimental::fixed_size_group<
-    PartitionSize, ParentGroup>> {
+template <size_t ChunkSize, typename ParentGroup>
+struct group_scope<sycl::ext::oneapi::experimental::chunk<
+    ChunkSize, ParentGroup>> {
   static constexpr __spv::Scope::Flag value = group_scope<ParentGroup>::value;
 };
 
@@ -174,15 +174,15 @@ bool GroupAll(ext::oneapi::experimental::ballot_group<ParentGroup> g,
     return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
   }
 }
-template <size_t PartitionSize, typename ParentGroup>
+template <size_t ChunkSize, typename ParentGroup>
 bool GroupAll(
-    ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>,
+    ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
     bool pred) {
   // GroupNonUniformAll doesn't support cluster size, so use a reduction
   return __spirv_GroupNonUniformBitwiseAnd(
       group_scope<ParentGroup>::value,
       static_cast<uint32_t>(__spv::GroupOperation::ClusteredReduce),
-      static_cast<uint32_t>(pred), PartitionSize);
+      static_cast<uint32_t>(pred), ChunkSize);
 }
 template <typename ParentGroup>
 bool GroupAll(ext::oneapi::experimental::tangle_group<ParentGroup>, bool pred) {
@@ -210,15 +210,15 @@ bool GroupAny(ext::oneapi::experimental::ballot_group<ParentGroup> g,
     return __spirv_GroupNonUniformAny(group_scope<ParentGroup>::value, pred);
   }
 }
-template <size_t PartitionSize, typename ParentGroup>
+template <size_t ChunkSize, typename ParentGroup>
 bool GroupAny(
-    ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>,
+    ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
     bool pred) {
   // GroupNonUniformAny doesn't support cluster size, so use a reduction
   return __spirv_GroupNonUniformBitwiseOr(
       group_scope<ParentGroup>::value,
       static_cast<uint32_t>(__spv::GroupOperation::ClusteredReduce),
-      static_cast<uint32_t>(pred), PartitionSize);
+      static_cast<uint32_t>(pred), ChunkSize);
 }
 template <typename ParentGroup>
 bool GroupAny(ext::oneapi::experimental::tangle_group<ParentGroup>, bool pred) {
@@ -327,12 +327,12 @@ GroupBroadcast(sycl::ext::oneapi::experimental::ballot_group<ParentGroup> g,
                                             WideOCLX, OCLId);
   }
 }
-template <size_t PartitionSize, typename ParentGroup, typename T, typename IdT>
+template <size_t ChunkSize, typename ParentGroup, typename T, typename IdT>
 EnableIfNativeBroadcast<T, IdT> GroupBroadcast(
-    ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup> g,
+    ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g,
     T x, IdT local_id) {
   // Remap local_id to its original numbering in ParentGroup
-  auto LocalId = g.get_group_linear_id() * PartitionSize + local_id;
+  auto LocalId = g.get_group_linear_id() * ChunkSize + local_id;
 
   // TODO: Refactor to avoid duplication after design settles.
   auto GroupLocalId = static_cast<typename GroupId<ParentGroup>::type>(LocalId);
@@ -341,9 +341,9 @@ EnableIfNativeBroadcast<T, IdT> GroupBroadcast(
   auto OCLId = detail::convertToOpenCLType(GroupLocalId);
 
   // NonUniformBroadcast requires Id to be dynamically uniform, which does not
-  // hold here; each partition is broadcasting a separate index. We could
+  // hold here; each chunk is broadcasting a separate index. We could
   // fallback to either NonUniformShuffle or a NonUniformBroadcast per
-  // partition, and it's unclear which will be faster in practice.
+  // chunk, and it's unclear which will be faster in practice.
   return __spirv_GroupNonUniformShuffle(group_scope<ParentGroup>::value,
                                         WideOCLX, OCLId);
 }
@@ -1298,12 +1298,10 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
     }                                                                          \
   }                                                                            \
                                                                                \
-  template <__spv::GroupOperation Op, size_t PartitionSize,                    \
+  template <__spv::GroupOperation Op, size_t ChunkSize,                        \
             typename ParentGroup, typename T>                                  \
   inline T Group##Instruction(                                                 \
-      ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>  \
-          g,                                                                   \
-      T x) {                                                                   \
+      ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g, T x)         \
     using ConvertedT = detail::ConvertToOpenCLType_t<T>;                       \
                                                                                \
     using OCLT = std::conditional_t<                                           \
@@ -1321,7 +1319,7 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
       constexpr auto OpInt =                                                   \
           static_cast<unsigned int>(__spv::GroupOperation::ClusteredReduce);   \
       return __spirv_GroupNonUniform##Instruction(Scope, OpInt, Arg,           \
-                                                  PartitionSize);              \
+                                                  ChunkSize);              \
     } else {                                                                   \
       T tmp;                                                                   \
       for (size_t Cluster = 0; Cluster < g.get_group_linear_range();           \

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

@@ -26,18 +26,18 @@ namespace sycl {
 inline namespace _V1 {
 namespace ext::oneapi::experimental {
 
-template <size_t PartitionSize, typename ParentGroup> class fixed_size_group;
+template <size_t ChunkSize, typename ParentGroup> class chunk;
 
-template <size_t PartitionSize, typename Group>
+template <size_t ChunkSize, typename Group>
 #ifdef __SYCL_DEVICE_ONLY__
 [[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_fixed_size_group)]]
 #endif
 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
                             std::is_same_v<Group, sycl::sub_group>,
-                        fixed_size_group<PartitionSize, Group>>
-get_fixed_size_group(Group group);
+                        chunk<ChunkSize, Group>>
+chunked_partition(Group group);
 
-template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
+template <size_t ChunkSize, typename ParentGroup> class chunk {
 public:
   using id_type = id<1>;
   using range_type = range<1>;
@@ -47,7 +47,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 
   id_type get_group_id() const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_SubgroupLocalInvocationId() / PartitionSize;
+    return __spirv_SubgroupLocalInvocationId() / ChunkSize;
 #else
     throw exception(make_error_code(errc::runtime),
                     "Non-uniform groups are not supported on host.");
@@ -56,7 +56,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 
   id_type get_local_id() const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_SubgroupLocalInvocationId() % PartitionSize;
+    return __spirv_SubgroupLocalInvocationId() % ChunkSize;
 #else
     throw exception(make_error_code(errc::runtime),
                     "Non-uniform groups are not supported on host.");
@@ -65,7 +65,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 
   range_type get_group_range() const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_SubgroupSize() / PartitionSize;
+    return __spirv_SubgroupSize() / ChunkSize;
 #else
     throw exception(make_error_code(errc::runtime),
                     "Non-uniform groups are not supported on host.");
@@ -74,7 +74,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 
   range_type get_local_range() const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return PartitionSize;
+    return ChunkSize;
 #else
     throw exception(make_error_code(errc::runtime),
                     "Non-uniform groups are not supported on host.");
@@ -137,57 +137,57 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
   fixed_size_group() {}
 #endif
 
-  friend fixed_size_group<PartitionSize, ParentGroup>
-  get_fixed_size_group<PartitionSize, ParentGroup>(ParentGroup g);
+  friend chunk<ChunkSize, ParentGroup>
+  chunked_partition<ChunkSize, ParentGroup>(ParentGroup g);
 
   friend sub_group_mask
-  sycl::detail::GetMask<fixed_size_group<PartitionSize, ParentGroup>>(
-      fixed_size_group<PartitionSize, ParentGroup> Group);
+  sycl::detail::GetMask<chunk<ChunkSize, ParentGroup>>(
+      chunk<ChunkSize, ParentGroup> Group);
 };
 
-template <size_t PartitionSize, typename Group>
+template <size_t ChunkSize, typename Group>
 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
                             std::is_same_v<Group, sycl::sub_group>,
-                        fixed_size_group<PartitionSize, Group>>
-get_fixed_size_group(Group group) {
+                        chunk<ChunkSize, Group>>
+chunked_partition(Group group) {
   (void)group;
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__NVPTX__)
   uint32_t loc_id = group.get_local_linear_id();
   uint32_t loc_size = group.get_local_linear_range();
-  uint32_t bits = PartitionSize == 32
+  uint32_t bits = ChunkSize == 32
                       ? 0xffffffff
-                      : ((1 << PartitionSize) - 1)
-                            << ((loc_id / PartitionSize) * PartitionSize);
+                      : ((1 << ChunkSize) - 1)
+                            << ((loc_id / ChunkSize) * ChunkSize);
 
-  return fixed_size_group<PartitionSize, sycl::sub_group>(
+  return chunk<ChunkSize, sycl::sub_group>(
       sycl::detail::Builder::createSubGroupMask<ext::oneapi::sub_group_mask>(
           bits, loc_size));
 #else
-  return fixed_size_group<PartitionSize, sycl::sub_group>();
+  return chunk<ChunkSize, sycl::sub_group>();
 #endif
 #else
   throw exception(make_error_code(errc::runtime),
                   "Non-uniform groups are not supported on host.");
 #endif
 }
 
-template <size_t PartitionSize, typename ParentGroup>
-struct is_user_constructed_group<fixed_size_group<PartitionSize, ParentGroup>>
+template <size_t ChunkSize, typename ParentGroup>
+struct is_user_constructed_group<chunk<ChunkSize, ParentGroup>>
     : std::true_type {};
 
 } // namespace ext::oneapi::experimental
 
 namespace detail {
-template <size_t PartitionSize, typename ParentGroup>
-struct is_fixed_size_group<
-    ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>>
+template <size_t ChunkSize, typename ParentGroup>
+struct is_chunk<
+    ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>>
     : std::true_type {};
 } // namespace detail
 
-template <size_t PartitionSize, typename ParentGroup>
+template <size_t ChunkSize, typename ParentGroup>
 struct is_group<
-    ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>>
+    ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>>
     : std::true_type {};
 
 } // namespace _V1

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

@@ -73,7 +73,7 @@ namespace ext::oneapi::experimental {
 
 // Forward declarations of non-uniform group types for algorithm definitions
 template <typename ParentGroup> class ballot_group;
-template <size_t PartitionSize, typename ParentGroup> class fixed_size_group;
+template <size_t ChunkSize, typename ParentGroup> class chunk;
 template <typename ParentGroup> class tangle_group;
 class opportunistic_group;
 

sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp → sycl/test-e2e/NonUniformGroups/chunk.cpp

@@ -13,19 +13,19 @@
 #include <vector>
 namespace syclex = sycl::ext::oneapi::experimental;
 
-template <size_t PartitionSize> class TestKernel;
+template <size_t ChunkSize> class TestKernel;
 
-template <size_t PartitionSize> void test() {
+template <size_t ChunkSize> void test() {
   sycl::queue Q;
 
   // Test for both the full sub-group size and a case with less work than a full
   // sub-group.
   for (size_t WGS : std::array<size_t, 2>{32, 16}) {
-    if (WGS < PartitionSize)
+    if (WGS < ChunkSize)
       continue;
 
     std::cout << "Testing for work size " << WGS << " and partition size "
-              << PartitionSize << std::endl;
+              << ChunkSize << std::endl;
 
     sycl::buffer<bool, 1> MatchBuf{sycl::range{WGS}};
     sycl::buffer<bool, 1> LeaderBuf{sycl::range{WGS}};
@@ -40,24 +40,24 @@ template <size_t PartitionSize> void test() {
             auto SG = item.get_sub_group();
             auto SGS = SG.get_local_linear_range();
 
-            auto Partition = syclex::get_fixed_size_group<PartitionSize>(SG);
+            auto Partition = syclex::chunked_partition<ChunkSize>(SG);
 
             bool Match = true;
-            Match &= (Partition.get_group_id() == (WI / PartitionSize));
-            Match &= (Partition.get_local_id() == (WI % PartitionSize));
-            Match &= (Partition.get_group_range() == (SGS / PartitionSize));
-            Match &= (Partition.get_local_range() == PartitionSize);
+            Match &= (Partition.get_group_id() == (WI / ChunkSize));
+            Match &= (Partition.get_local_id() == (WI % ChunkSize));
+            Match &= (Partition.get_group_range() == (SGS / ChunkSize));
+            Match &= (Partition.get_local_range() == ChunkSize);
             MatchAcc[WI] = Match;
             LeaderAcc[WI] = Partition.leader();
           };
-      CGH.parallel_for<TestKernel<PartitionSize>>(NDR, KernelFunc);
+      CGH.parallel_for<TestKernel<ChunkSize>>(NDR, KernelFunc);
     });
 
     sycl::host_accessor MatchAcc{MatchBuf, sycl::read_only};
     sycl::host_accessor LeaderAcc{LeaderBuf, sycl::read_only};
     for (int WI = 0; WI < WGS; ++WI) {
       assert(MatchAcc[WI] == true);
-      assert(LeaderAcc[WI] == ((WI % PartitionSize) == 0));
+      assert(LeaderAcc[WI] == ((WI % ChunkSize) == 0));
     }
   }
 }