[SYCL] Added ext_oneapi_non_uniform_groups aspect (#10902)

The errors that you get if you try to e.g. create a `ballot_group` in
e.g. HIP backend are not instructive. This PR introduces
`ext_oneapi_non_uniform_groups` and uses the
`[[__sycl_detail__::__uses_aspects__(aspect::foo)]]` annotations
detailed in:


https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md#changes-to-dpc-headers

to give a clear runtime error in the case that a programmer tries to
execute a kernel that uses non-uniform groups on a device that does not
supported the non-uniform group extension.

---------

Signed-off-by: Jack Kirk <[email protected]>
Co-authored-by: Greg Lueck <[email protected]>

Author: JackAKirk

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -62,6 +62,7 @@ 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">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -110,7 +111,7 @@ 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_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd, AspectExt_oneapi_non_uniform_groups],
     []>;
 // 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

@@ -37,14 +37,10 @@ https://github.com/intel/llvm/issues
 
 == Dependencies
 
-This extension is written against the SYCL 2020 revision 6 specification.  All
+This extension is written against the SYCL 2020 revision 7 specification.  All
 references below to the "core SYCL specification" or to section numbers in the
 SYCL specification refer to that revision.
 
-This extension also depends on the following other SYCL extensions:
-
-* link:https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc[sycl_ext_oneapi_root_group]
-
 
 == Status
 
@@ -55,6 +51,15 @@ incompatible ways before it is finalized.  *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
+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
+kernel is submitted to the queue.
+
 == Overview
 
 Many modern hardware architectures support flexible sub-divisions of
@@ -73,8 +78,9 @@ needed in function documentation.
 
 NOTE: The first version of this extension only supports partitioning of
 sub-groups. It is expected that in the future, these functions will be expanded
-to also allow partitioning of root-groups, work-groups and user-constructed
-groups.
+to also allow partitioning of
+link:https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc[root-groups],
+work-groups and user-constructed groups.
 
 
 == Specification
@@ -99,6 +105,21 @@ implementation supports.
  feature-test macro always has this value.
 |===
 
+=== Extension to `enum class aspect`
+
+[source]
+----
+namespace sycl {
+enum class aspect {
+  ...
+  ext_oneapi_non_uniform_groups
+}
+}
+----
+
+If a SYCL device has the `ext_oneapi_non_uniform_groups` aspect,
+then it supports the non-uniform groups described in the next sections.
+
 === Control Flow
 
 The SYCL specification defines
@@ -130,7 +151,7 @@ model topology used by SYCL kernels. These groups are implicitly created by an
 implementation when a SYCL kernel function is enqueued. The following group
 types are fixed topology groups:
 
-- `root_group` (if sycl_ext_oneapi_root_group is supported)
+- `root_group` (if link:https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc[sycl_ext_oneapi_root_group] is supported)
 - `group`
 - `sub_group`
 

sycl/include/sycl/device_aspect_macros.hpp

@@ -278,6 +278,11 @@
 #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
+#endif
+
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
 // __SYCL_ASPECT(host, 0)
 #define __SYCL_ANY_DEVICE_HAS_host__ 0
@@ -547,3 +552,8 @@
 //__SYCL_ASPECT(ext_intel_esimd, 53)
 #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
+#endif

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

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/aspects.hpp>
 #include <sycl/detail/pi.h>            // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
 #include <sycl/exception.hpp>          // for runtime_error
@@ -27,6 +28,10 @@ namespace ext::oneapi::experimental {
 template <typename ParentGroup> class ballot_group;
 
 template <typename Group>
+#ifdef __SYCL_DEVICE_ONLY__
+[[__sycl_detail__::__uses_aspects__(
+    sycl::aspect::ext_oneapi_non_uniform_groups)]]
+#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>>
@@ -142,6 +147,7 @@ inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
 get_ballot_group(Group group, bool predicate) {
   (void)group;
 #ifdef __SYCL_DEVICE_ONLY__
+#if defined(__SPIR__) || defined(__NVPTX__)
   // ballot_group partitions into two groups using the predicate
   // Membership mask for one group is negation of the other
   sub_group_mask mask = sycl::ext::oneapi::group_ballot(group, predicate);
@@ -150,6 +156,7 @@ get_ballot_group(Group group, bool predicate) {
   } else {
     return ballot_group<sycl::sub_group>(~mask, predicate);
   }
+#endif
 #else
   (void)predicate;
   throw runtime_error("Non-uniform groups are not supported on host device.",

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

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/aspects.hpp>
 #include <sycl/detail/pi.h>            // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/type_traits.hpp> // for is_fixed_size_group, is_group
 #include <sycl/exception.hpp>          // for runtime_error
@@ -27,6 +28,10 @@ namespace ext::oneapi::experimental {
 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)]]
+#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>>

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

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/aspects.hpp>
 #include <sycl/detail/pi.h>                   // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/type_traits.hpp>        // for is_group, is_user_cons...
 #include <sycl/exception.hpp>                 // for runtime_error
@@ -26,8 +27,13 @@ namespace ext::oneapi::experimental {
 class opportunistic_group;
 
 namespace this_kernel {
-inline opportunistic_group get_opportunistic_group();
-}
+#ifdef __SYCL_DEVICE_ONLY__
+[[__sycl_detail__::__uses_aspects__(
+    sycl::aspect::ext_oneapi_non_uniform_groups)]]
+#endif
+inline opportunistic_group
+get_opportunistic_group();
+} // namespace this_kernel
 
 class opportunistic_group {
 public:

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

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/aspects.hpp>
 #include <sycl/detail/pi.h>                   // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/type_traits.hpp>        // for is_group, is_user_cons...
 #include <sycl/exception.hpp>                 // for runtime_error
@@ -26,6 +27,10 @@ namespace ext::oneapi::experimental {
 template <typename ParentGroup> class tangle_group;
 
 template <typename Group>
+#ifdef __SYCL_DEVICE_ONLY__
+[[__sycl_detail__::__uses_aspects__(
+    sycl::aspect::ext_oneapi_non_uniform_groups)]]
+#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>>
@@ -148,6 +153,8 @@ get_tangle_group(Group group) {
   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.");
 #endif
 #else
   throw runtime_error("Non-uniform groups are not supported on host device.",

sycl/include/sycl/info/aspects.def

@@ -48,3 +48,4 @@ __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/source/detail/device_impl.cpp

@@ -554,6 +554,11 @@ bool device_impl::has(aspect Aspect) const {
             &support, nullptr) == PI_SUCCESS;
     return call_successful && support;
   }
+  case aspect::ext_oneapi_non_uniform_groups: {
+    return (this->getBackend() == backend::ext_oneapi_level_zero) ||
+           (this->getBackend() == backend::opencl) ||
+           (this->getBackend() == backend::ext_oneapi_cuda);
+  }
   }
   throw runtime_error("This device aspect has not been implemented yet.",
                       PI_ERROR_INVALID_DEVICE);

sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp

@@ -2,6 +2,7 @@
 // RUN: %{run} %t.out
 
 // REQUIRES: gpu
+// UNSUPPORTED: hip
 // REQUIRES: sg-32
 
 #include <sycl/sycl.hpp>