[SYCL] Remove get_child_group() (#13482)

get_child_group() was originally designed assuming there would be a
fixed
hierarchy of groups, and is therefore incompatible with extensions that
would
modify the hierarchy (e.g., "clusters" on NVIDIA GPUs).

This interface should be replaced by a more general and explicit
mechanism for
partitioning groups that does not make such assumptions.

Signed-off-by: John Pennycook <[email protected]>

---------

Signed-off-by: John Pennycook <[email protected]>

Author: Pennycook

sycl/doc/extensions/experimental/sycl_ext_oneapi_root_group.asciidoc

@@ -20,7 +20,7 @@
 == Notice
 
 [%hardbreaks]
-Copyright (C) 2022-2023 Intel Corporation.  All rights reserved.
+Copyright (C) 2022-2024 Intel Corporation.  All rights reserved.
 
 Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
 of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
@@ -407,36 +407,7 @@ _Returns_: The number of work-groups or work-items of `Scope` hierarchy level wi
 
 `root_group` provides an alternative representation of the work-items executing
 an ND-range kernel and exposes equivalent functionality to `sycl::nd_item` for
-querying a work-item's position in the global range. In order to provide access
-to information pertaining to a work-item's position in the work-group or
-sub-group local range, `root_group` needs to provide a new mechanism to access
-instances of the `sycl::group` and `sycl::sub_group` classes. The
-`get_child_group` function provides a general form of this mechanism, allowing
-developers to move down the hierarchy of fixed topology groups.
-
-[source,c++]
-----
-template <typename Group>
-/* type of child group */ get_child_group(Group g);
-----
-_Constraints_: `Group` must be one of `root_group` or `sycl::group`.
-
-_Returns_: An instance of another fixed topology group type, representing the
-child of group _g_ to which the calling work-item belongs. If `Group` is
-`root_group`, the child group type is `sycl::group` with the same
-dimensionality. If `Group` is `sycl::group`, the child group type is
-`sycl::sub_group`.
-
-NOTE: Although `sycl::sub_group` is a fixed topology group, it is currently
-the lowest level of the hierarchy and cannot be passed to `get_child_group`.
-
-NOTE: This extension does not provide a `get_parent_group` function because it
-would be easy to use incorrectly. It is good practice for a function accepting
-a group _g_ to only use the work-items in that group, to assist developers in
-reasoning about requirements of the call context (e.g. converged control flow).
-Dividing a group into its children is consistent with this practice, whereas
-accessing the parent group is not. Developers seeking this functionality should
-use the free function queries instead.
+querying a work-item's position in the global range.
 
 
 === Synchronizing a `root_group`

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

@@ -79,16 +79,6 @@ template <int Dimensions> class root_group {
   sycl::nd_item<Dimensions> it;
 };
 
-template <int Dimensions>
-group<Dimensions> get_child_group(root_group<Dimensions> g) {
-  (void)g;
-  return this_group<Dimensions>();
-}
-
-template <int Dimensions> sycl::sub_group get_child_group(group<Dimensions> g) {
-  (void)g;
-  return this_sub_group();
-}
 namespace this_work_item {
 template <int Dimensions> root_group<Dimensions> get_root_group() {
   return sycl::ext::oneapi::this_work_item::get_nd_item<Dimensions>()
@@ -118,7 +108,7 @@ void group_barrier(ext::oneapi::experimental::root_group<dimensions> G,
   // workaround that's not intended to reduce the bar for SPIR-V modules
   // acceptance, but rather make a pessimistic case work until we have full
   // support for the device barrier built-in from backends.
-  const auto ChildGroup = ext::oneapi::experimental::get_child_group(G);
+  const auto ChildGroup = ext::oneapi::experimental::this_group<dimensions>();
   if (ChildGroup.get_group_linear_range() == 1) {
     group_barrier(ChildGroup);
   } else {

sycl/test-e2e/GroupAlgorithm/root_group.cpp

@@ -80,7 +80,7 @@ void testRootGroupFunctions() {
   const auto props = sycl::ext::oneapi::experimental::properties{
       sycl::ext::oneapi::experimental::use_root_sync};
 
-  constexpr int testCount = 10;
+  constexpr int testCount = 9;
   sycl::buffer<bool> testResultsBuf{sycl::range{testCount}};
   const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
   q.submit([&](sycl::handler &h) {
@@ -103,16 +103,6 @@ void testRootGroupFunctions() {
                 root.get_local_linear_id() == root.get_local_id().get(0);
             testResults[7] = root.get_group_linear_range() == 1;
             testResults[8] = root.get_local_linear_range() == WorkGroupSize;
-
-            const auto child =
-                sycl::ext::oneapi::experimental::get_child_group(root);
-            const auto grandchild =
-                sycl::ext::oneapi::experimental::get_child_group(child);
-            testResults[9] = child == it.get_group();
-            static_assert(
-                std::is_same_v<std::remove_cv<decltype(grandchild)>::type,
-                               sycl::sub_group>,
-                "get_child_group(sycl::group) must return a sycl::sub_group");
           }
         });
   });