Add sycl_khr_group_interface extension

[Pennycook]

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.

[TApplencourt]

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 just ndit.khr_get_group()?)
• khr::group_interface::get_item?

[Pennycook]

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 just ndit.khr_get_group()?)
• khr::group_interface::get_item?

I don't think so. When vendors introduce extensions they use sycl::ext::vendor::, not sycl::ext::vendor::extension_name.

The way I see it, although putting something in the khr:: 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.

[keryell]

That looks good!

[keryell]

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

[Pennycook]

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

Good idea. I've added rank, rank_dynamic and static_extent in 008ac23.

[AerialMantis]

One thing we didn't discuss -- that only became apparent to me after making the change -- is whether people are still happy with the name get_item. We could now alternatively call the sycl::member_item constructor (but this wouldn't work with ADL), or could rename the function something like get_member or get_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 to get_member_item, it would avoid any confusion with the current item 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.

[Pennycook]

I was thinking about this when I was reviewing this again just now, we could potentially change this to get_member_item, it would avoid any confusion with the current item 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]

WG discussed, and converged on get_member_item.

[Pennycook]

CTS tests are implemented in KhronosGroup/SYCL-CTS#1024.

The DPC++ implementation that passes the CTS tests is under review in intel/llvm#17595.

[gmlueck]

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

[PeterTh]

I've prototyped 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 of using id_type = id<Dimensions>;, and the same for range (see here).

[Pennycook]

I've prototyped 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 of using id_type = id<Dimensions>;, and the same for range (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.

[tomdeakin]

WG approved to merge

[gmlueck]

Sorry for the late comment. Our recent discussion about passing queue by value vs. by reference got me thinking about similar issues with this KHR.

[gmlueck]

Should this instead be:

Suggested change

	member_item<Group> get_member_item(Group g) noexcept;
	member_item<Group> get_member_item(const Group& g) noexcept;

Likewise for all the other functions in this KHR that take a Group, group, or sub_group parameter by value.

[Pennycook]

I'm not opposed to this direction (and I've made the change in 985a67a), 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 both Group g and const 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 take Group and the same functions in SYCL-Next take const 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, because const Group& is less restrictive than Group (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?

[gmlueck]

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.

[AerialMantis]

Looks good! I still have a couple of comments open, but neither of them are blockers to merging the PR.

[gmlueck]

WG approved to merge

@tomdeakin: I did not merge this because 985a67a 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]

WG approved to merge (again)