-
Notifications
You must be signed in to change notification settings - Fork 808
[SYCL][L0 v2] Fix bindless array image subregion copies #20952
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 our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: sycl
Are you sure you want to change the base?
Conversation
4c2a346 to
186e439
Compare
186e439 to
89881bd
Compare
…dation + UR region mapping)
89881bd to
17d7288
Compare
9932f88 to
f0f42ea
Compare
f0f42ea to
8581964
Compare
| } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D || | ||
| ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) { | ||
| Region->depth = 1; | ||
| // Runtime validation of Origin values based on image type |
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 have to admit, I'm not super well versed in image support in L0. Can you describe what the change is, and why it fixes the problem?
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.
The code is simplified now. Still, basically the change in level_zero adapter code fixes mapping of image region dimensions which is different in UR and L0 for 1DARRAY images. In particular, numbers of layers to copy in UR: Region-depth, L0: Height. Starting index UR: Origin->z, L0: OriginY
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.
Also 1D images need adjusting - UR allows Height and Depth to be 0 (only Width set to image size), while L0 needs them to be 1.
AlexeySachkov
left a comment
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.
Thank you for fixing this!
Conceptually this looks good to me (although I'm not a code owner anymore), but I do have a few notes/questions
| SrcImgDesc.depth}; | ||
| sycl::range<3> DestImageSize = {DestImgDesc.width, DestImgDesc.height, | ||
| DestImgDesc.depth}; | ||
| // If this is a multi-layer array image, use the layer count; otherwise, use |
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 see that L0 does not support arrays of 3d images, but briefly looking at the extension specification I couldn't find any restriction about that on the SYCL level.
Did I just missed it, or is it indeed missing from the spec?
If it is missing, I'm not asking to add it in this PR, but creating an issue to remind us about this would be helpful
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.
SYCL also does not support arrays of 3D images. The specification is not written clearly, but sentences like "... the 3rd dimension of the ranges represent the arrays' layer(s) being copied, regardless of whether the copy is performed on a 1D or 2D image array." indicate this. You can also see it in the code:
| // Image arrays must only be 1D or 2D. |
| (ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY && Origin->y == 0) || | ||
| (ZeImageDesc.type == ZE_IMAGE_TYPE_2D && Origin->z == 0) || | ||
| (ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) || |
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.
Do we also need to assert y == 0 for 1D array and z == 0 for 2D array?
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 believe the current code is correct as-is. The assertions properly validate that:
- Unused dimensions are zero (y and z for 1D, z for 2D)
- Array layer dimensions (z for both 1D and 2D arrays) are allowed to be non-zero
Fix bindless array image subregion copies on Level Zero v2 (SYCL validation + UR region mapping)
Enable tests (previously failing and disabled with no tracking):
bindless_images/array/read_write_1d_subregion.cpp
bindless_images/array/read_write_2d_subregion.cpp