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 10 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 @@ -67,6 +67,7 @@ def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
def AspectUsm_system_allocator : Aspect<"usm_system_allocator">;
def AspectUsm_restricted_shared_allocations : Aspect<"usm_restricted_shared_allocations">;
def AspectHost : Aspect<"host">;
def AspectExt_oneapi_non_uniform_groups : Aspect<"ext_oneapi_non_uniform_groups">;
defvar AllUSMAspects = [AspectUsm_device_allocations, AspectUsm_host_allocations,
AspectUsm_shared_allocations, AspectUsm_system_allocations, AspectUsm_atomic_host_allocations,
AspectUsm_atomic_shared_allocations];
Expand Down Expand Up @@ -109,7 +110,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_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, 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
28 changes: 20 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 Down Expand Up @@ -73,8 +69,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 +96,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 +142,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 @@ -273,6 +273,11 @@
#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_level_reference__ 0
#endif

#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_non_uniform_groups__
// __SYCL_ASPECT(ext_oneapi_non_uniform_groups, 53)
#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 @@ -537,3 +542,8 @@
//__SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52)
#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_level_reference__ 0
#endif

#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_non_uniform_groups__
// __SYCL_ASPECT(ext_oneapi_non_uniform_groups, 53)
#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_non_uniform_groups__ 0
#endif
9 changes: 9 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 @@ -30,7 +31,13 @@ template <typename Group>
inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
std::is_same_v<Group, sycl::sub_group>,
ballot_group<Group>>
#ifdef __SYCL_DEVICE_ONLY__
get_ballot_group [[__sycl_detail__::__uses_aspects__(
sycl::aspect::ext_oneapi_non_uniform_groups)]] (Group group,
Copy link
Contributor

Choose a reason for hiding this comment

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

Is this the right location to place an attribute for a function declaration? I thought the attribute went before the return type.

Copy link
Contributor Author

@JackAKirk JackAKirk Sep 8, 2023 •
edited
Loading

Choose a reason for hiding this comment

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

Yeah the docs suggest this, but if I put it before the return type : i.e before std::enable_if_t.. I get:

error: '__uses_aspects__' attribute cannot be applied to types
   33 | [[__sycl_detail__::__uses_aspects__(

If I put it before the inline too then it works, but I can't have any variation like this that doesn't then give a warning when building the compiler like:

warning: extra tokens at end of #endif directive

I have tried putting it in a bunch of places, and the only way I can compile and run without any warnings/errors and get it working correctly is like how I've done it in the PR!

I could make it superficially different so it also avoids the #else statement but it looks a mess, so I decided this way was best.

Copy link
Contributor

Choose a reason for hiding this comment

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

According to cppreference, the attribute is allowed in either place: before the inline or directly after the function name. So, I think what you have is fine.

(https://en.cppreference.com/w/cpp/language/attributes)

FWIW, I presume you could do this if you wanted to place the attribute before the inline:

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>>
get_ballot_group(Group group, bool predicate);

Copy link
Contributor Author

Choose a reason for hiding this comment

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

You're right this does work, I thought I tried it and it didn't. It is better like that. I'll change it.

bool predicate);
#else
get_ballot_group(Group group, bool predicate);
#endif

template <typename ParentGroup> class ballot_group {
public:
Expand Down Expand Up @@ -142,6 +149,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 +158,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
6 changes: 6 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 @@ -30,7 +31,12 @@ template <size_t PartitionSize, typename Group>
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>>
#ifdef __SYCL_DEVICE_ONLY__
get_fixed_size_group [[__sycl_detail__::__uses_aspects__(
sycl::aspect::ext_oneapi_non_uniform_groups)]] (Group group);
#else
get_fixed_size_group(Group group);
#endif

template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
public:
Expand Down
9 changes: 8 additions & 1 deletion 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,14 @@ namespace ext::oneapi::experimental {
class opportunistic_group;

namespace this_kernel {
#ifdef __SYCL_DEVICE_ONLY__
inline opportunistic_group get_opportunistic_group
[[__sycl_detail__::__uses_aspects__(
sycl::aspect::ext_oneapi_non_uniform_groups)]] ();
#else
inline opportunistic_group get_opportunistic_group();
}
#endif
} // namespace this_kernel

class opportunistic_group {
public:
Expand Down
8 changes: 8 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 @@ -29,7 +30,12 @@ template <typename Group>
inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
std::is_same_v<Group, sycl::sub_group>,
tangle_group<Group>>
#ifdef __SYCL_DEVICE_ONLY__
get_tangle_group [[__sycl_detail__::__uses_aspects__(
sycl::aspect::ext_oneapi_non_uniform_groups)]] (Group group);
#else
get_tangle_group(Group group);
#endif

template <typename ParentGroup> class tangle_group {
public:
Expand Down Expand Up @@ -148,6 +154,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 @@ -47,3 +47,4 @@ __SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49)
__SYCL_ASPECT(ext_oneapi_mipmap, 50)
__SYCL_ASPECT(ext_oneapi_mipmap_anisotropy, 51)
__SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52)
__SYCL_ASPECT(ext_oneapi_non_uniform_groups, 53)
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 @@ -546,6 +546,11 @@ bool device_impl::has(aspect Aspect) const {
sizeof(pi_bool), &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

#include <sycl/sycl.hpp>
#include <vector>
Expand Down