Movatterモバイル変換


[0]ホーム

URL:


Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly

Sign up
Appearance settings

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

Merged
gmlueck merged 28 commits intoKhronosGroup:mainfromPennycook:khr_group_interface
May 29, 2025

Conversation

Pennycook
Copy link
Contributor

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.

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.
khr::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
Copy link
Contributor

TApplencourt commentedOct 7, 2024
edited
Loading

I like the usage of namespace.

    khr::work_group<1> g = ndit.get_group();    khr::work_item it = get_item(g);

Few questions (sorry, trying to understand this new khr mechanism), should it be :

  • khr::group_interface::work_item
  • ndit.khr_group_interface_get_group? (or justndit.khr_get_group()?)
  • khr::group_interface::get_item?

@Pennycook
Copy link
ContributorAuthor

I like the usage of namespace.

    khr::work_group<1> g = ndit.get_group();    khr::work_item it = get_item(g);

Few questions (sorry, trying to understand this new khr mechanism), should it be :

  • khr::group_interface::work_item
  • ndit.khr_group_interface_get_group? (or justndit.khr_get_group()?)
  • khr::group_interface::get_item?

I don't think so. When vendors introduce extensions they usesycl::ext::vendor::, notsycl::ext::vendor::extension_name.

The way I see it, although putting something in thekhr:: namespace reserves a name, the SYCL WG is in complete control over any names that are introduced there. Future extensions might have to be more creative with their naming, but there's no risk of surprise conflicts.

TApplencourt reacted with thumbs up emoji

Copy link
Member

@keryellkeryell left a comment

Choose a reason for hiding this comment

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

That looks good!

@keryell
Copy link
Member

Adding the 2 rank functions frommdspan too?https://gitlab.khronos.org/sycl/Specification/-/issues/630

rank(), rank_dynamic() and static_extent() are all exposed as staticmember functions in mdspan, so we follow that precedent here.
@Pennycook
Copy link
ContributorAuthor

Adding the 2 rank functions frommdspan too?https://gitlab.khronos.org/sycl/Specification/-/issues/630

Good idea. I've addedrank,rank_dynamic andstatic_extent in008ac23.

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).
Copy link
Collaborator

@AerialMantisAerialMantis left a 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 = /* ... */;
Copy link
Collaborator

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?

Copy link
ContributorAuthor

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.

Copy link
Collaborator

@AerialMantisAerialMantisMay 15, 2025
edited
Loading

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.

Copy link
ContributorAuthor

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.

AerialMantis reacted with thumbs up emoji
member_item<Group> get_item(Group g) noexcept;
----

_Constraints_: `Group` is `work_group` or `sub_group`.
Copy link
Collaborator

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.

Copy link
ContributorAuthor

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.

Copy link
Collaborator

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.

Copy link
ContributorAuthor

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.

AerialMantis reacted with thumbs up emoji
@AerialMantis
Copy link
Collaborator

One thing we didn't discuss -- that only became apparent to me after making the change -- is whether people are still happy with the nameget_item. We could now alternatively call thesycl::member_item constructor (but this wouldn't work with ADL), or could rename the function something likeget_member orget_member_item (but both of these are more verbose).

I was thinking about this when I was reviewing this again just now, we could potentially change this toget_member_item, it would avoid any confusion with the currentitem class but it's a bit more wordy. I don't have a strong opinion on this either, but maybe something to discuss with the rest of the working group.

@PennycookPennycook added the AgendaTo be discussed during a SYCL committee meeting labelFeb 28, 2025
@Pennycook
Copy link
ContributorAuthor

I was thinking about this when I was reviewing this again just now, we could potentially change this toget_member_item, it would avoid any confusion with the currentitem class but it's a bit more wordy. I don't have a strong opinion on this either, but maybe something to discuss with the rest of the working group.

I've requested to put this on the agenda for next week's meeting.

@tomdeakin
Copy link
Contributor

WG discussed, and converged onget_member_item.

@PennycookPennycook removed the AgendaTo be discussed during a SYCL committee meeting labelMar 7, 2025
Intended to avoid confusion between sycl::item and khr::member_item.
@Pennycook
Copy link
ContributorAuthor

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.
@gmlueck
Copy link
Contributor

@illuhad : I think we are waiting for your input on this. Do you agree that this is implementable?

@PeterTh
Copy link

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.using id_type = sycl::id<Dimensions>; instead ofusing id_type = id<Dimensions>;, and the same forrange (see here).

@Pennycook
Copy link
ContributorAuthor

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).

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...

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.using id_type = sycl::id<Dimensions>; instead ofusing id_type = id<Dimensions>;, and the same forrange (see here).

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.

@PennycookPennycook added the AgendaTo be discussed during a SYCL committee meeting labelApr 30, 2025
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.
@PennycookPennycook removed the AgendaTo be discussed during a SYCL committee meeting labelMay 15, 2025
@tomdeakin
Copy link
Contributor

WG approved to merge

Copy link
Contributor

@gmlueckgmlueck left a 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;
Copy link
Contributor

Choose a reason for hiding this comment

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

Should this instead be:

Suggested change
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.

Copy link
ContributorAuthor

@PennycookPennycookMay 16, 2025
edited
Loading

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?

Copy link
Contributor

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.

Pennycook reacted with thumbs up emoji
Copy link
Collaborator

@AerialMantisAerialMantis left a 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.

@gmlueck
Copy link
Contributor

WG approved to merge

@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.

@gmlueck
Copy link
Contributor

WG approved to merge (again)

@gmlueckgmlueck merged commitd903d2f intoKhronosGroup:mainMay 29, 2025
3 checks passed
@PennycookPennycook deleted the khr_group_interface branchJune 3, 2025 08:04
Sign up for freeto join this conversation on GitHub. Already have an account?Sign in to comment
Reviewers

@keryellkeryellkeryell approved these changes

@gmlueckgmlueckgmlueck approved these changes

@TApplencourtTApplencourtTApplencourt approved these changes

@AerialMantisAerialMantisAerialMantis approved these changes

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.

7 participants
@Pennycook@TApplencourt@keryell@AerialMantis@tomdeakin@gmlueck@PeterTh

[8]ページ先頭

©2009-2025 Movatter.jp