diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index cbfc1074d29fd..ca940bddf01c7 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/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. diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc index 8a8b8af19a958..d168684af72ee 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc +++ b/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` diff --git a/sycl/include/sycl/device_aspect_macros.hpp b/sycl/include/sycl/device_aspect_macros.hpp index fbc79f1235c05..1b1c3f8bc301d 100644 --- a/sycl/include/sycl/device_aspect_macros.hpp +++ b/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 diff --git a/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp index f20f2f12fb8b9..567438f3fc837 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp @@ -8,6 +8,7 @@ #pragma once +#include #include // for PI_ERROR_INVALID_DEVICE #include // for is_group, is_user_cons... #include // for runtime_error @@ -27,6 +28,10 @@ namespace ext::oneapi::experimental { template class ballot_group; template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::__uses_aspects__( + sycl::aspect::ext_oneapi_non_uniform_groups)]] +#endif inline std::enable_if_t> && std::is_same_v, ballot_group> @@ -142,6 +147,7 @@ inline std::enable_if_t> && 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(~mask, predicate); } +#endif #else (void)predicate; throw runtime_error("Non-uniform groups are not supported on host device.", diff --git a/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp index 2fbbda9b45669..7e130a7b78032 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp @@ -8,6 +8,7 @@ #pragma once +#include #include // for PI_ERROR_INVALID_DEVICE #include // for is_fixed_size_group, is_group #include // for runtime_error @@ -27,6 +28,10 @@ namespace ext::oneapi::experimental { template class fixed_size_group; template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::__uses_aspects__( + sycl::aspect::ext_oneapi_non_uniform_groups)]] +#endif inline std::enable_if_t> && std::is_same_v, fixed_size_group> diff --git a/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp index 224cffc8bf77d..2d5cca911314c 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp @@ -8,6 +8,7 @@ #pragma once +#include #include // for PI_ERROR_INVALID_DEVICE #include // for is_group, is_user_cons... #include // 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: diff --git a/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp index 3d21f2bfd2268..a8a0a6bf3102e 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp @@ -8,6 +8,7 @@ #pragma once +#include #include // for PI_ERROR_INVALID_DEVICE #include // for is_group, is_user_cons... #include // for runtime_error @@ -26,6 +27,10 @@ namespace ext::oneapi::experimental { template class tangle_group; template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::__uses_aspects__( + sycl::aspect::ext_oneapi_non_uniform_groups)]] +#endif inline std::enable_if_t> && std::is_same_v, tangle_group> @@ -148,6 +153,8 @@ get_tangle_group(Group group) { return tangle_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.", diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 3f82b93a885ca..b28a4b0be2f15 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/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) diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index b48f044471d22..5de2a19c7728d 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/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); diff --git a/sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp b/sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp index 1542824adde1e..2f043c5bed711 100644 --- a/sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp +++ b/sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp @@ -2,6 +2,7 @@ // RUN: %{run} %t.out // REQUIRES: gpu +// UNSUPPORTED: hip // REQUIRES: sg-32 #include