[SYCL] Restore the nd_range default constructor - #22949
Conversation
PR intel#22908 removed nd_range's default constructor to align with the SYCL 2020 specification, which does not define one. That broke code which relies on default-constructing an nd_range and assigning it later, most notably CuTe's launch_policy used by PyTorch's XPU flash-attention kernels. Restore the constructor as a backwards-compatibility extension and add a test that checks it zero-initializes the constituent ranges. Fixes intel#22940
|
Note for reviewers - please do not merge it right now. We need a few days to discuss this and decide if we want to put it to SYCL 2020 (most likely) or to create a oneapi extension. |
Please, keep me updated on this. |
|
Not sure how long I'll take to update the spec. I've added macros to keep this until the next major release. I suggest to proceed that way, hope it will be added to the spec soon, so I'll remove the macros. P.S. now it should look like we didn't remove this, but just guarded to be removed prior to next major release. |
|
@KseniyaTikhomirova could you take a look please? This change is desired by sycl-tla team. |
| nd_range<Dimensions> &operator=(const nd_range<Dimensions> &rhs) = default; | ||
| nd_range<Dimensions> &operator=(nd_range<Dimensions> &&rhs) = default; | ||
|
|
||
| #ifndef __INTEL_PREVIEW_BREAKING_CHANGES |
There was a problem hiding this comment.
On pytorch side we can't accept this solution in the release version of the compiler, neither closed source DPC++ nor open source dpclang. We need this resolved before the next release. Can you educate me what exactly is an issue to step out from the spec on this matter and continue to have nd_range default compiler defined?
There was a problem hiding this comment.
@dvrogozh you mean you can't accept this macro? By the next release you mean 7.3.0? Because 7.2.0 will contain this ctor. I believe we'll resolve this by 7.3.0, this is just to unblock folks that need this right now.
Can you educate me what exactly is an issue to step out from the spec on this matter and continue to have nd_range default compiler defined?
@gmlueck do you think we can restore this right now without any guards? Note - we had it for a long time, but I removed it recently, so it won't be like adding a new API but rather restoring the existed one.
There was a problem hiding this comment.
Pytorch stopped using preview macro long time ago and it was an effort dedicated to that due to Pytorch being production environment. And we don't want to introduce it again.
this is just to unblock folks that need this right now.
This solution is OK as intermediate engineering solution, but that's showstopper from my perspective for the next release which seems to be 7.3.0 per your comment.
There was a problem hiding this comment.
This ctor is added under "if not defined INTEL_PREVIEW...", so it'll be available even in 7.3.0, but might be removed prior to 8.0.0. I believe we'll add it to SYCL 2020 or create a oneapi extension by this time.
Does it work for you?
There was a problem hiding this comment.
oh, ok, I misread the condition looking into the PR early morning. Then reformulating the concern, this PR is fine right now but gives a risk to the next major release 8.0.0 if you will decide to drop this ctor.
There was a problem hiding this comment.
I am not opposed to this change, but it might make sense to wait until later this week if there is no hurry. I'm hoping to bring KhronosGroup/SYCL-Docs#1043 up for discussing on Thursday of this week. If the Khronos WG decides to add the default constructor as an "errata" fix, then we won't need this PR at all. We can just keep the default constructor as it is now.
There was a problem hiding this comment.
Without this change intel/llvm pytorch CI is blocked. Longer it's broken more risk for other issues to appear unnoticed contributing to the triage bucket.
There was a problem hiding this comment.
Maybe I'm reading this PR wrong. I thought it was just removing the default constructor in -fpreview-breaking-changes mode? There's no change at all in the regular mode, right?
@dvrogozh, I think you said above that PyTorch doesn't use the -fpreview-breaking-changes mode. If my understanding about this PR is correct, I don't think it will have any effect on PyTorch at all.
There was a problem hiding this comment.
I am talking about pytorch CI defined in this intel/llvm repository. You have a workflow which validates most recent builds of the dpclang compiler with the pytorch. See https://github.com/intel/llvm/blob/sycl/.github/workflows/sycl-pytorch-build-and-test.yml.
There was a problem hiding this comment.
I understand now. Sorry, I didn't read the PR description closely enough. I see now that this is undoing the recent change in #22908, which removed that default constructor.
Feel free to merge this PR now if you want. If the WG decides later to add the default constructor, we can make another PR to remove the #ifndef.
|
Hi @arnavprabhu, |
|
Hi @KornevNikita, I’ve updated my GitHub email settings accordingly. Please let me know if there’s anything else I need to address on the PR. Thanks, |
5 participants
Fixes #22940
Problem
PR #22908 removed
nd_range's default constructor (nd_range() = default;)to align the class with the SYCL 2020 specification, which does not declare
one. That broke code which default-constructs an
nd_rangeand assigns itlater. The most visible casualty is CuTe's
launch_policy(
include/cute/util/compat/launch_policy.hpp), used by PyTorch's XPUflash-attention kernels, which holds an
nd_range<3>member that getsdefault-initialized:
Change
Restore
nd_range() = default;. All three members (globalSize,localSize,offset) arerange/idobjects, both of which keep zero-initializingdefault constructors, so a default-constructed
nd_rangeis well-defined andproduces zeroed ranges and offset.
Test plan
sycl/test/basic_tests/nd_range.cppto default-construct annd_range<3>and assert the constituent ranges and offset arezero-initialized.