[SYCL][Docs] Move sycl_ext_oneapi_non_uniform_groups to experimental (#11535)

This commit makes the following changes to allow us to consider the
sycl_ext_oneapi_non_uniform_groups extensions as experimentally
supported:
* Make the extension aspects per-group to allow the CUDA backend to not
support tangle_group yet.
* Define the SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS feature test macro.
* Move the extension document to experimental.

---------

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -62,7 +62,10 @@ def AspectExt_oneapi_mipmap : Aspect<"ext_oneapi_mipmap">;
 def AspectExt_oneapi_mipmap_anisotropy : Aspect<"ext_oneapi_mipmap_anisotropy">;
 def AspectExt_oneapi_mipmap_level_reference : Aspect<"ext_oneapi_mipmap_level_reference">;
 def AspectExt_intel_esimd : Aspect<"ext_intel_esimd">;
-def AspectExt_oneapi_non_uniform_groups : Aspect<"ext_oneapi_non_uniform_groups">;
+def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group">;
+def AspectExt_oneapi_fixed_size_group : Aspect<"ext_oneapi_fixed_size_group">;
+def AspectExt_oneapi_opportunistic_group : Aspect<"ext_oneapi_opportunistic_group">;
+def AspectExt_oneapi_tangle_group : Aspect<"ext_oneapi_tangle_group">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -111,7 +114,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_bindless_images_shared_usm, AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm,
     AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_memory_export,
     AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_interop_semaphore_export,
-    AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd, AspectExt_oneapi_non_uniform_groups],
+    AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd,
+    AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group, AspectExt_oneapi_tangle_group],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc

@@ -44,17 +44,18 @@ SYCL specification refer to that revision.
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 
 == Backend support status
 
-The APIs in this extension may be used only on a device that has
-`aspect::ext_oneapi_non_uniform_groups`. The application must check that the
+The APIs in this extension may be used only on a device that has one or more of
+the xref:ext-aspects[extension aspects]. The application must check that the
 device has this aspect before submitting a kernel using any of the APIs in this
 extension. If the application fails to do this, the implementation throws a
 synchronous exception with the `errc::kernel_not_supported` error code when the
@@ -105,20 +106,33 @@ implementation supports.
  feature-test macro always has this value.
 |===
 
+[#ext-aspects]
 === Extension to `enum class aspect`
 
 [source]
 ----
 namespace sycl {
 enum class aspect {
   ...
-  ext_oneapi_non_uniform_groups
+  ext_oneapi_ballot_group
+  ext_oneapi_fixed_size_group
+  ext_oneapi_opportunistic_group
+  ext_oneapi_tangle_group
 }
 }
 ----
 
-If a SYCL device has the `ext_oneapi_non_uniform_groups` aspect,
-then it supports the non-uniform groups described in the next sections.
+If a SYCL device has these aspects, that device supports the non-uniform groups
+as follows:
+
+[%header,cols="2,3"]
+|===
+| Aspect | Supported group
+| `ext_oneapi_ballot_group` | `ballot_group`
+| `ext_oneapi_fixed_size_group` | `fixed_size_group`
+| `ext_oneapi_opportunistic_group` | `opportunistic_group`
+| `ext_oneapi_tangle_group` | `tangle_group`
+|===
 
 === Control Flow
 

sycl/include/sycl/device_aspect_macros.hpp

@@ -278,9 +278,24 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_intel_esimd__ 0
 #endif
 
-#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_non_uniform_groups__
-// __SYCL_ASPECT(ext_oneapi_non_uniform_groups, 54)
-#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_non_uniform_groups__ 0
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_ballot_group__
+// __SYCL_ASPECT(ext_oneapi_ballot_group, 54)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_ballot_group__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_fixed_size_group__
+// __SYCL_ASPECT(ext_oneapi_fixed_size_group, 55)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_fixed_size_group__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_opportunistic_group__
+// __SYCL_ASPECT(ext_oneapi_opportunistic_group, 56)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_opportunistic_group__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_tangle_group__
+// __SYCL_ASPECT(ext_oneapi_tangle_group, 57)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_tangle_group__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
@@ -553,7 +568,22 @@
 #define __SYCL_ANY_DEVICE_HAS_ext_intel_esimd__ 0
 #endif
 
-#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_non_uniform_groups__
-// __SYCL_ASPECT(ext_oneapi_non_uniform_groups, 54)
-#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_non_uniform_groups__ 0
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_ballot_group__
+// __SYCL_ASPECT(ext_oneapi_ballot_group, 54)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_ballot_group__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_fixed_size_group__
+// __SYCL_ASPECT(ext_oneapi_fixed_size_group, 55)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_fixed_size_group__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_opportunistic_group__
+// __SYCL_ASPECT(ext_oneapi_opportunistic_group, 56)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_opportunistic_group__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_tangle_group__
+// __SYCL_ASPECT(ext_oneapi_tangle_group, 57)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_tangle_group__ 0
 #endif

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

@@ -29,13 +29,12 @@ template <typename ParentGroup> class ballot_group;
 
 template <typename Group>
 #ifdef __SYCL_DEVICE_ONLY__
-[[__sycl_detail__::__uses_aspects__(
-    sycl::aspect::ext_oneapi_non_uniform_groups)]]
+[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_ballot_group)]]
 #endif
 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
                             std::is_same_v<Group, sycl::sub_group>,
-                        ballot_group<Group>>
-get_ballot_group(Group group, bool predicate);
+                        ballot_group<Group>> get_ballot_group(Group group,
+                                                              bool predicate);
 
 template <typename ParentGroup> class ballot_group {
 public:

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

@@ -29,8 +29,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group;
 
 template <size_t PartitionSize, typename Group>
 #ifdef __SYCL_DEVICE_ONLY__
-[[__sycl_detail__::__uses_aspects__(
-    sycl::aspect::ext_oneapi_non_uniform_groups)]]
+[[__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>,

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

@@ -29,10 +29,9 @@ class opportunistic_group;
 namespace this_kernel {
 #ifdef __SYCL_DEVICE_ONLY__
 [[__sycl_detail__::__uses_aspects__(
-    sycl::aspect::ext_oneapi_non_uniform_groups)]]
+    sycl::aspect::ext_oneapi_opportunistic_group)]]
 #endif
-inline opportunistic_group
-get_opportunistic_group();
+inline opportunistic_group get_opportunistic_group();
 } // namespace this_kernel
 
 class opportunistic_group {

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

@@ -28,13 +28,11 @@ template <typename ParentGroup> class tangle_group;
 
 template <typename Group>
 #ifdef __SYCL_DEVICE_ONLY__
-[[__sycl_detail__::__uses_aspects__(
-    sycl::aspect::ext_oneapi_non_uniform_groups)]]
+[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_tangle_group)]]
 #endif
 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
                             std::is_same_v<Group, sycl::sub_group>,
-                        tangle_group<Group>>
-get_tangle_group(Group group);
+                        tangle_group<Group>> get_tangle_group(Group group);
 
 template <typename ParentGroup> class tangle_group {
 public:
@@ -152,9 +150,10 @@ get_tangle_group(Group group) {
   sub_group_mask mask = sycl::ext::oneapi::group_ballot(group, true);
   return tangle_group<sycl::sub_group>(mask);
 #elif defined(__NVPTX__)
-  // TODO: Construct from compiler-generated mask
-  static_assert(false,
-                "tangle_group is not currently supported on this platform.");
+  // TODO: Construct from compiler-generated mask. Return an invalid group in
+  //       in the meantime. CUDA devices will report false for the tangle_group
+  //       support aspect so kernels launch should ensure this is never run.
+  return tangle_group<sycl::sub_group>(0);
 #endif
 #else
   throw runtime_error("Non-uniform groups are not supported on host device.",

sycl/include/sycl/info/aspects.def

@@ -48,4 +48,7 @@ __SYCL_ASPECT(ext_oneapi_mipmap, 50)
 __SYCL_ASPECT(ext_oneapi_mipmap_anisotropy, 51)
 __SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52)
 __SYCL_ASPECT(ext_intel_esimd, 53)
-__SYCL_ASPECT(ext_oneapi_non_uniform_groups, 54)
+__SYCL_ASPECT(ext_oneapi_ballot_group, 54)
+__SYCL_ASPECT(ext_oneapi_fixed_size_group, 55)
+__SYCL_ASPECT(ext_oneapi_opportunistic_group, 56)
+__SYCL_ASPECT(ext_oneapi_tangle_group, 57)

sycl/source/detail/device_impl.cpp

@@ -554,11 +554,19 @@ bool device_impl::has(aspect Aspect) const {
             &support, nullptr) == PI_SUCCESS;
     return call_successful && support;
   }
-  case aspect::ext_oneapi_non_uniform_groups: {
+  case aspect::ext_oneapi_ballot_group:
+  case aspect::ext_oneapi_fixed_size_group:
+  case aspect::ext_oneapi_opportunistic_group: {
     return (this->getBackend() == backend::ext_oneapi_level_zero) ||
            (this->getBackend() == backend::opencl) ||
            (this->getBackend() == backend::ext_oneapi_cuda);
   }
+  case aspect::ext_oneapi_tangle_group: {
+    // TODO: tangle_group is not currently supported for CUDA devices. Add when
+    //       implemented.
+    return (this->getBackend() == backend::ext_oneapi_level_zero) ||
+           (this->getBackend() == backend::opencl);
+  }
   }
   throw runtime_error("This device aspect has not been implemented yet.",
                       PI_ERROR_INVALID_DEVICE);

sycl/source/feature_test.hpp.in

@@ -96,6 +96,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_PREFETCH 1
 #define SYCL_EXT_INTEL_CACHE_CONTROLS 1
 #define SYCL_EXT_INTEL_FP_CONTROL 1
+#define SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS 1
 
 #ifndef __has_include
 #define __has_include(x) 0