- Notifications
You must be signed in to change notification settings - Fork71
Add sycl_khr_group_interface extension#638
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.
By clicking “Sign up for GitHub”, you agree to ourterms of service andprivacy statement. We’ll occasionally send you account related emails.
Already on GitHub?Sign in to your account
Uh oh!
There was an error while loading.Please reload this page.
Conversation
This extension introduces an alternative interface for groups ofwork-items, offering several improvements over the SYCL 2020 interface:- Shorter names for member functions, dropping the get_ prefix.- Cleaner separation between properties of a group (e.g., a group id) and properties of the calling work-item (e.g., its id within a group).- Clearer distinction between "group" concept and "work_group" class.- New work_item class to represent a single work-item within a specific parent group. This class also satisfies the group concept, modeling a group containing a single work-item.
dc92d73
to93e784e
Comparekhr::work_item is not templated on dimensions.
The synopsis previously assumed that the extents would be onedimensional, but it must match the ParentGroup. There is no suitableshorthand alias or exposition-only description for this case definedby mdspan, so replace the type with a comment.
TApplencourt commentedOct 7, 2024 • edited
Loading Uh oh!
There was an error while loading.Please reload this page.
edited
Uh oh!
There was an error while loading.Please reload this page.
I like the usage of namespace.
Few questions (sorry, trying to understand this new khr mechanism), should it be :
|
I don't think so. When vendors introduce extensions they use The way I see it, although putting something in the |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
That looks good!
Adding the 2 rank functions from |
rank(), rank_dynamic() and static_extent() are all exposed as staticmember functions in mdspan, so we follow that precedent here.
Good idea. I've added |
25b35fe
tod59ccb9
CompareUh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
The name __group__ is consistent with the names of the swizzle classes.
Co-authored-by: Greg Lueck <gregory.m.lueck@intel.com>
Uh oh!
There was an error while loading.Please reload this page.
The previous "not user constructible" wording was based on the description ofexisting group types. The new wording explains why khr::group, khr::sub_group,and khr::work_item have limited constructors and also points to functions thatcan be used to acquire instances of these types.
The group interface (which we hope to replace with a concept) should notdictate whether or not a function is constexpr unless all groups must have thesame behavior.Since we have two existing groups (work-group and sub-group) where it isimpossible to determine their sizes at compile-time, it would be misleading todeclare functions relating to their size as constexpr.Making this change to the interface requires the member functions to beduplicated for every class. I've used this as an opportunity to clarify thebehavior of these functions in a way that is specific to each class (e.g.,highlighting that a sub-group's indices are with respect to its parentwork-group).
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Looks great! Just a few minor suggestions.
using range_type = /* ... */; | ||
using extents_type = /* ... */; // C++23 | ||
using size_type = /* ... */; | ||
static constexpr int dimensions = /* ... */; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
If we're looking to align with ISO C++ should we consider calling thisrank
instead?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
There is astatic constexpr rank()
function that's available starting with C++23.
dimensions
is already part of the SYCL 2020 group interface (in the sense that it's defined bygroup
andsub_group
).
I assumed we wouldn't want to remove that any time soon, so made it part of this new interface as well.
AerialMantisMay 15, 2025 • edited
Loading Uh oh!
There was an error while loading.Please reload this page.
edited
Uh oh!
There was an error while loading.Please reload this page.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
That's true, though we are adding additional interfaces to the group classes, I suspect with the intention of removing the old interfaces at some point. We could do this fordimensions
as well, have an additional value rank which is equal todimensions
, and then we could look to deprecatedimensions
in the future.
Though saying that, this is not really an issue, just that it would be nice to align with ISO C++ naming, so it may not be worth it.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Good point, I hadn't considered thatdimensions
andrank()
are returning the same thing here!
I think deprecatingdimensions
in the future would make sense. But we can't introduce arank
value in its place, because then we'd have a variable calledrank
and a function calledrank()
in the same scope.rank()
is how it's spelt inmdspan
andextents
, so once the minimum version of SYCL hits C++23 we'll be perfectly aligned.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
member_item<Group> get_item(Group g) noexcept; | ||
---- | ||
_Constraints_: `Group` is `work_group` or `sub_group`. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
I wonder if we should base this constrain on a type trait so that we can support additional group types withget_item
in the future without having to modify the definition of this API.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
I agree we should move towards a trait (or concept) eventually, but I didn't want to hold up this initial KHR with trying to define which traits we need. The IWOCL presentation and subsequent discussion suggested that we might need multiple traits (e.g., to distinguish between groups that can/cannot synchronize, groups that can/cannot be partitioned, etc).
Because of the way extensions work, I envisioned that any extension that introduces a new group and wants to be compatible with this interface could just introduce a new overload ofget_item()
, rather than us having to revisit this extension and update an existing trait. Doing things this way also ensures that we don't make a function available by mistake, which I fear could be very easy if we introduce a trait too early.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
So I agree we will likely end up requiring multiple traits, so we may not want this to be a generic group trait, as we may introduce a group type in the future which we don't want to supportget_item
for.
Though I'm a little concerned about this approach, as it effectively means introducing new overloads for new group types which would increase the complexity, for example if we introduceget_item
forwork_group
andsub_group
now and then later introduceroot_group
,cluster_group
, etc, we would end up with something like:
template <typename Group, typename = std::enable_if_t<std::is_same_v<Group, work_group> || std::is_same_v<Group, sub_group>>::type>void get_item(Group g) noexcept {}void get_item(root_group g) noexcept {}void get_item(cluster_group g) noexcept {}
Where as if we had a trait in this extension, something likeis_item_indexable
(not a great name but I'm sure we could think of something better), then we could support additional types by simply making that group type conform tois_item_indexable
. I feel like this would be cleaner.
template <typename Group, typename = std::enable_if_t<is_item_indexable<Group>>::type>void get_item(Group g) noexcept {}
Saying that I don't want to hold up this PR, so perhaps we could consider this as a follow up change.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
We're on the same page. The original presentation that proposed that we should have a group concept calledindexable_item_group
, sois_item_indexable
is aligned with that proposal.
I think we should defer this conversation until a later date, though, perhaps even when we discuss the possibility of this feature becoming core. Whether we'd encode this as a concept or a set of traits will be dependent on whether we can assume C++20 as a minimum version, and that will depend on which version of SYCL this intercepts.
I also want to point out that the specification is deliberately vague about how this should be implemented, saying only:
_Constraints_: [code]#Group# is [code]#work_group# or [code]#sub_group#.
If we don't like the idea ofroot_group
adding a new overload ofget_item
, we could just add a relaxation that says something to the effect of: "Ifroot_group
is available, theConstraints ofget_member_item
are modified to permitroot_group
". So I think we can achieve what you want to without exposing a user-facing trait.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
I was thinking about this when I was reviewing this again just now, we could potentially change this to |
I've requested to put this on the agenda for next week's meeting. |
WG discussed, and converged on |
Intended to avoid confusion between sycl::item and khr::member_item.
CTS tests are implemented inKhronosGroup/SYCL-CTS#1024. The DPC++ implementation that passes the CTS tests is under review inintel/llvm#17595. |
The change indc1c0f5 missed a few instancesof constexpr, resulting in inconsistencies between header and functionsynopses. This commit fixes that such that size() is only constexpr formember_item.
@illuhad : I think we are waiting for your input on this. Do you agree that this is implementable? |
PeterTh commentedApr 30, 2025
I'veprototyped this extension in SimSYCL and I've noticed some most likely unintended behaviour in GCC -- see this minimal example:https://godbolt.org/z/zosojE1zG Curiously, Clang and MSVC completely ignore this issue. I wasn't able to find exactly what GCC is basing this decision on in a quick search of the current C++ standard, but on the other hand, I think it's pretty clear that this isn't great (and even if it was, GCC is an important compiler target). As such, I think it might be better if the relevant parts of the code in the spec were adjusted to nip this problem in the bud. This just requires a pretty simple spelling change of types in 4 lines of code, i.e. |
Interesting find. This is weird! After a little bit of searching I eventually found this:https://bugs.llvm.org/show_bug.cgi?id=15559 -- the text in the specification has changed so I can't find this exactly, but it looks like both GCC and Clang are correct...
I'm going to stick this on the agenda for Thursday's WG meeting. Not because I'm opposed to the proposed fix, but because I want to make sure there isn't another solution before I make the change. It's also not clear to me whether there's a precedent for something like this already in the specification, or whether it's an implementation issue rather than a specification one. |
Within the scope of a group definition, sycl::id<> and sycl::range<> are types,whereas id() and range() are functions. Some compilers issue a diagnostic whenthe same name is used like this, but others do not; adding sycl:: ensures thatthe header is valid code across all compilers.
WG approved to merge |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Sorry for the late comment. Our recent discussion about passingqueue
by value vs. by reference got me thinking about similar issues with this KHR.
}; | ||
template <typename Group> | ||
member_item<Group> get_member_item(Group g) noexcept; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Should this instead be:
member_item<Group> get_member_item(Group g) noexcept; | |
member_item<Group> get_member_item(constGroup& g) noexcept; |
Likewise for all the other functions in this KHR that take aGroup
,group
, orsub_group
parameter by value.
PennycookMay 16, 2025 • edited
Loading Uh oh!
There was an error while loading.Please reload this page.
edited
Uh oh!
There was an error while loading.Please reload this page.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
I'm not opposed to this direction (and I've made the change in985a67a), but one thing we should give some thought to is if we can also fix the existing group algorithms that take a group by value.
Having overloads for bothGroup g
andconst Group& g
defined for the same function would lead to ambiguity, so I think the only way to fix this would be to say that the functions in SYCL 2020 takeGroup
and the same functions in SYCL-Next takeconst Group&
. It wouldn't make sense to deprecate these functions in between, because we're not actually removing anything.
I think all implementations of the group algorithms are header-only because they accept user-defined types, so this shouldn't be an ABI break concern. I also think that the vast majority of user would continue to work, becauseconst Group&
is less restrictive thanGroup
(but I'm not prepared to rule out that there might be some weird corner cases that require some adjustment).
Could we make this change in SYCL-Next, or are we stuck?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Yes, I think we need to reexamine all the APIs in the core SYCL spec that take objects by value. It probably does make sense to change them in SYCL-Next even if it does cause some weird corner cases to break. If there are cases like this, users will need to adjust in order to compile applications in SYCL-Next mode.
I'm much more concerned about retaining ABI compatibility between SYCL 2020 and SYCL-Next. It would be a major pain if vendors needed to ship two versions of every library that uses SYCL.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Looks good! I still have a couple of comments open, but neither of them are blockers to merging the PR.
@tomdeakin: I did not merge this because985a67a was made after the WG approved. If you would like this merged prior to the next WG meeting, please add another comment saying so. Otherwise, we can call for it to be merged again at the next meeting. |
WG approved to merge (again) |
d903d2f
intoKhronosGroup:mainUh oh!
There was an error while loading.Please reload this page.
This extension introduces an alternative interface for groups of work-items, offering several improvements over the SYCL 2020 interface:
Shorter names for member functions, dropping the get_ prefix.
Cleaner separation between properties of a group (e.g., a group id) and properties of the calling work-item (e.g., its id within a group).
Clearer distinction between "group" concept and "work_group" class.
New work_item class to represent a single work-item within a specific parent group. This class also satisfies the group concept, modeling a group containing a single work-item.