[SYCL] Update root_group extension to use this_work_item namespace

sycl/doc/extensions/experimental/sycl_ext_oneapi_root_group.asciidoc

@@ -471,18 +471,15 @@ class nd_item {
    sycl::ext::oneapi::experimental::root_group<Dimensions> ext_oneapi_get_root_group() const;
 };
 
-namespace ext {
-namespace oneapi {
-namespace experimental {
-namespace this_kernel {
+namespace ext::oneapi::experimental {
+namespace this_work_item {
 
 template <int Dimensions>
 root_group<Dimensions> get_root_group();
 
-} // namespace this_kernel
-} // namespace experimental
-} // namespace oneapi
-} // namespace ext
+}
+
+} // namespace ext::oneapi::experimental
 } // namespace sycl
 ----
 
@@ -505,6 +502,31 @@ a `sycl::nd_range` argument.
 _Returns_: A `root_group` instance representing the root-group to which the
 calling work-item belongs.
 
+=== Deprecated functionality
+
+The functionality in this section was previously part of this extension, but is
+now deprecated.
+
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
+namespace this_kernel {
+
+template <int Dimensions>
+root_group<Dimensions> get_root_group();
+
+} // namespace this_kernel
+
+}
+----
+
+[source,c++]
+----
+template <int Dimensions>
+root_group<Dimensions> get_root_group();
+----
+_Effects_: Equivalent to `return this_work_item::get_root_group()`.
 
 == Implementation notes
 

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

@@ -89,12 +89,21 @@ template <int Dimensions> sycl::sub_group get_child_group(group<Dimensions> g) {
   (void)g;
   return this_sub_group();
 }
-
-namespace this_kernel {
+namespace this_work_item {
 template <int Dimensions> root_group<Dimensions> get_root_group() {
   return sycl::ext::oneapi::this_work_item::get_nd_item<Dimensions>()
       .ext_oneapi_get_root_group();
 }
+} // namespace this_work_item
+
+namespace this_kernel {
+template <int Dimensions>
+__SYCL_DEPRECATED(
+    "use sycl::ext::oneapi::experimental::this_work_item::get_root_group() "
+    "instead")
+root_group<Dimensions> get_root_group() {
+  this_work_item::get_root_group<Dimensions>();
+}
 } // namespace this_kernel
 
 } // namespace ext::oneapi::experimental

sycl/test-e2e/GroupAlgorithm/root_group.cpp

@@ -52,7 +52,8 @@ void testRootGroup() {
           sycl::group_barrier(root);
 
           root =
-              sycl::ext::oneapi::experimental::this_kernel::get_root_group<1>();
+              sycl::ext::oneapi::experimental::this_work_item::get_root_group<
+                  1>();
           int sum = data[root.get_local_id()] +
                     data[root.get_local_range() - root.get_local_id() - 1];
           sycl::group_barrier(root);

sycl/test/warnings/free_functions_deprecation.cpp

@@ -16,5 +16,8 @@ int main() {
   // expected-warning@+1 {{'this_id<1>' is deprecated: use nd_range kernel and sycl::ext::oneapi::this_work_item::get_nd_item() instead}}
   (void)sycl_exp::this_id<1>();
 
+  // expected-warning@+1 {{'get_root_group<1>' is deprecated: use sycl::ext::oneapi::experimental::this_work_item::get_root_group() instead}}
+  (void)sycl_exp::this_kernel::get_root_group<1>();
+
   return 0;
 }