Skip to content

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

Sign up for GitHub

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

Open
mateuszpn wants to merge 16 commits into
base: sycl
Choose a base branch
from
+37 −36
Open

[SYCL][L0 v2] Fix bindless array image subregion copies #20952

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
16 commits
Select commit Hold shift + click to select a range
17d7288
Fix bindless array image subregion copies on Level Zero v2 (SYCL vali…
mateuszpn Dec 22, 2025
8581964
fix getImageRegionHelper() in l0 adapter
mateuszpn Dec 22, 2025
7b5df3a
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Jan 5, 2026
39e810f
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Jan 15, 2026
f50bfbf
Fix 1D array image region mapping
mateuszpn Jan 19, 2026
2908463
comment added
mateuszpn Jan 19, 2026
5eab493
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Jan 20, 2026
34dc5a8
fix for 1D images
mateuszpn Jan 20, 2026
d6c2865
Merge branch 'subregion-tests' of github.com:mateuszpn/llvm into subr…
mateuszpn Jan 20, 2026
3aac708
fix normalization code
mateuszpn Jan 20, 2026
95e4c25
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Jan 21, 2026
e76d246
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Jan 23, 2026
c0fae93
apply comment
mateuszpn Jan 23, 2026
e0ae69e
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Jan 26, 2026
ecf4d1f
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Jan 29, 2026
baa84e4
Merge branch 'intel:sycl' into subregion-tests
mateuszpn Feb 2, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 8 additions & 4 deletions sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -204,10 +204,14 @@ verify_sub_copy(const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
static_cast<bool>(result[2]));
};

sycl::range<3> SrcImageSize = {SrcImgDesc.width, SrcImgDesc.height,
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
Copy link
Contributor

@AlexeySachkov AlexeySachkov Jan 23, 2026

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

Copy link
Contributor Author

@mateuszpn mateuszpn Jan 23, 2026

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:

llvm/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp

Line 157 in 92aea26

// Image arrays must only be 1D or 2D.

// the depth dimension (following the logic in fill_image_type() ).
sycl::range<3> SrcImageSize = {
SrcImgDesc.width, SrcImgDesc.height,
SrcImgDesc.array_size > 1 ? SrcImgDesc.array_size : SrcImgDesc.depth};
sycl::range<3> DestImageSize = {
DestImgDesc.width, DestImgDesc.height,
DestImgDesc.array_size > 1 ? DestImgDesc.array_size : DestImgDesc.depth};

if (isOutOfRange(SrcImageSize, SrcOffset, CopyExtent) ||
isOutOfRange(DestImageSize, DestOffset, CopyExtent)) {
Expand Down
3 changes: 0 additions & 3 deletions sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,9 +1,6 @@
// REQUIRES: aspect-ext_oneapi_bindless_images
// REQUIRES: aspect-ext_oneapi_image_array

// UNSUPPORTED: level_zero
// UNSUPPORTED-INTENDED: Undetermined issue causing enqueue process to fail.

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

Expand Down
4 changes: 0 additions & 4 deletions sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,10 +1,6 @@
// REQUIRES: aspect-ext_oneapi_bindless_images
// REQUIRES: aspect-ext_oneapi_image_array

// UNSUPPORTED: level_zero
// UNSUPPORTED-INTENDED: Undetermined issue causing data and invalid pointer
// errors.

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

Expand Down
54 changes: 29 additions & 25 deletions unified-runtime/source/adapters/level_zero/image_common.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -731,47 +731,51 @@ ur_result_t getImageRegionHelper(ze_image_desc_t ZeImageDesc,
UR_ASSERT(Origin, UR_RESULT_ERROR_INVALID_VALUE);
UR_ASSERT(Region, UR_RESULT_ERROR_INVALID_VALUE);

if (ZeImageDesc.type == ZE_IMAGE_TYPE_1D ||
ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY) {
Region->height = 1;
Region->depth = 1;
} else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D ||
ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) {
Region->depth = 1;
}

#ifndef NDEBUG
// Validate Origin constraints based on image type
UR_ASSERT((ZeImageDesc.type == ZE_IMAGE_TYPE_1D && Origin->y == 0 &&
Origin->z == 0) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY && Origin->z == 0) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY && Origin->y == 0) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_2D && Origin->z == 0) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) ||
Comment on lines +738 to 740
Copy link
Contributor

@AlexeySachkov AlexeySachkov Jan 23, 2026

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?

Copy link
Contributor Author

@mateuszpn mateuszpn Jan 23, 2026

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

(ZeImageDesc.type == ZE_IMAGE_TYPE_3D),
UR_RESULT_ERROR_INVALID_VALUE);

UR_ASSERT(Region->width && Region->height && Region->depth,
// Validate Region width is non-zero
UR_ASSERT(Region->width != 0, UR_RESULT_ERROR_INVALID_VALUE);

// Validate Region depth for 1D arrays contains layer count
UR_ASSERT(ZeImageDesc.type != ZE_IMAGE_TYPE_1DARRAY || Region->depth != 0,
UR_RESULT_ERROR_INVALID_VALUE);

// Validate Region depth for 2D arrays contains layer count
UR_ASSERT(ZeImageDesc.type != ZE_IMAGE_TYPE_2DARRAY || Region->depth != 0,
UR_RESULT_ERROR_INVALID_VALUE);
UR_ASSERT(
(ZeImageDesc.type == ZE_IMAGE_TYPE_1D && Region->height == 1 &&
Region->depth == 1) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY && Region->depth == 1) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_2D && Region->depth == 1) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) ||
(ZeImageDesc.type == ZE_IMAGE_TYPE_3D),
UR_RESULT_ERROR_INVALID_VALUE);
#endif // !NDEBUG

uint32_t OriginX = ur_cast<uint32_t>(Origin->x);
uint32_t OriginY = ur_cast<uint32_t>(Origin->y);
uint32_t OriginZ = ur_cast<uint32_t>(Origin->z);

uint32_t Width = ur_cast<uint32_t>(Region->width);
uint32_t Height = (ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY)
? ZeImageDesc.arraylevels
: ur_cast<uint32_t>(Region->height);
uint32_t Depth = (ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY)
? ZeImageDesc.arraylevels
: ur_cast<uint32_t>(Region->depth);
uint32_t Height = ur_cast<uint32_t>(Region->height);
uint32_t Depth = ur_cast<uint32_t>(Region->depth);

// Normalize Region dimensions based on image type
if (ZeImageDesc.type == ZE_IMAGE_TYPE_1D) {
// 1D images: height and depth must be 1
Height = 1;
Depth = 1;
} else if (ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY) {
// UR uses depth for 1D array layers, but Level Zero uses height
OriginY = OriginZ;
OriginZ = 0;
Height = Depth;
Depth = 1;
} else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D) {
// 2D images: depth must be 1
Depth = 1;
}

ZeRegion = {OriginX, OriginY, OriginZ, Width, Height, Depth};

Expand Down