Skip to content

[SYCL] Added ext_oneapi_non_uniform_groups aspect #10902

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 20 commits into from
Sep 26, 2023
Merged

[SYCL] Added ext_oneapi_non_uniform_groups aspect #10902

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
75fbe0f
Added asserts to non-supporting platforms
Aug 21, 2023
fe5ab18
Lint.
Aug 21, 2023
32961b5
Lint part 2.
Aug 21, 2023
d885fdd
Impl ext_oneapi_non_uniform_groups aspect.
Aug 30, 2023
06206f8
Lint.
Aug 30, 2023
a59229b
Adjust mention of root_groups in doc.
Aug 31, 2023
ee5e1c2
Fix comma.
Aug 31, 2023
a8e85aa
Added ext_oneapi_non_uniform_groups aspect to __TestAspectList.
Aug 31, 2023
9ea80ca
Place new aspect before "]"
Aug 31, 2023
fd2f62f
Marked hip unsupported in fixed_size_group.cpp
Aug 31, 2023
f75ab44
Added `Backend support status` section to doc.
Sep 8, 2023
79df2ab
Moved __uses_aspects__attribute before type.
Sep 8, 2023
558c8d2
Update sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_group…
JackAKirk Sep 8, 2023
b34be0e
Update sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_group…
JackAKirk Sep 8, 2023
9c1f3a7
Format tangle_group attr.
Sep 8, 2023
9546f8d
Format.
Sep 8, 2023
be657c9
Moved aspect definition to appropriate place.
Sep 18, 2023
03ca377
Merge branch 'sycl' into non-uniform-hip-errors
JackAKirk Sep 20, 2023
2be3514
Fix last merge commit.
Sep 20, 2023
07fda3d
Small aesthetic non functional change.
Sep 22, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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">;
Expand Down Expand Up @@ -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.
Expand Down
37 changes: 29 additions & 8 deletions sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand All @@ -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
Expand All @@ -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
Expand All @@ -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
Expand Down Expand Up @@ -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`

Expand Down
10 changes: 10 additions & 0 deletions sycl/include/sycl/device_aspect_macros.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
7 changes: 7 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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>>
Expand Down Expand Up @@ -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);
Expand All @@ -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.",
Expand Down
5 changes: 5 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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>>
Expand Down
10 changes: 8 additions & 2 deletions sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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:
Expand Down
7 changes: 7 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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>>
Expand Down Expand Up @@ -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.",
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/info/aspects.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
5 changes: 5 additions & 0 deletions sycl/source/detail/device_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) ||
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@Pennycook I also wanted to check that this correctly covers all the SPIRV cases: Are

(this->getBackend() == backend::ext_oneapi_level_zero) ||
           (this->getBackend() == backend::opencl)

The only two backends supporting the spirv non uniform groups?

Thanks

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think these are the only two that we've tested, at least. I'm happy to leave this as it is for now -- if we can prove to ourselves later that things should run elsewhere, we can update the aspect.

(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);
Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/NonUniformGroups/fixed_size_group.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,7 @@
// RUN: %{run} %t.out

// REQUIRES: gpu
// UNSUPPORTED: hip
// REQUIRES: sg-32

#include <sycl/sycl.hpp>
Expand Down