From 17d7288b397b5ed2373a905cd05be903556fd70b Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Mon, 22 Dec 2025 13:21:06 +0000 Subject: [PATCH 1/2] Fix bindless array image subregion copies on Level Zero v2 (SYCL validation + UR region mapping) --- sycl/source/handler.cpp | 10 +++-- .../array/read_write_1d_subregion.cpp | 3 -- .../array/read_write_2d_subregion.cpp | 4 -- .../adapters/level_zero/image_common.cpp | 41 ++++++++++++------- 4 files changed, 33 insertions(+), 25 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index b03f7f0cfd855..3fefe89771976 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -210,10 +210,12 @@ verify_sub_copy(const ext::oneapi::experimental::image_descriptor &SrcImgDesc, static_cast(result[2])); }; - sycl::range<3> SrcImageSize = {SrcImgDesc.width, SrcImgDesc.height, - SrcImgDesc.depth}; - sycl::range<3> DestImageSize = {DestImgDesc.width, DestImgDesc.height, - DestImgDesc.depth}; + 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)) { diff --git a/sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp b/sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp index baedbb11c3c77..40cde9309b091 100644 --- a/sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp +++ b/sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp @@ -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 diff --git a/sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp b/sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp index 5c86922c950a7..1abeb19a35cd4 100644 --- a/sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp +++ b/sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp @@ -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 diff --git a/unified-runtime/source/adapters/level_zero/image_common.cpp b/unified-runtime/source/adapters/level_zero/image_common.cpp index 35dbc739e2d28..6d786061ead37 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/image_common.cpp @@ -731,19 +731,17 @@ 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) { + if (ZeImageDesc.type == ZE_IMAGE_TYPE_1D) { Region->height = 1; Region->depth = 1; - } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D || - ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) { + } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D) { Region->depth = 1; } #ifndef NDEBUG 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) || (ZeImageDesc.type == ZE_IMAGE_TYPE_3D), @@ -754,7 +752,7 @@ ur_result_t getImageRegionHelper(ze_image_desc_t ZeImageDesc, 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_1DARRAY && Region->height == 1) || (ZeImageDesc.type == ZE_IMAGE_TYPE_2D && Region->depth == 1) || (ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) || (ZeImageDesc.type == ZE_IMAGE_TYPE_3D), @@ -766,12 +764,25 @@ ur_result_t getImageRegionHelper(ze_image_desc_t ZeImageDesc, uint32_t OriginZ = ur_cast(Origin->z); uint32_t Width = ur_cast(Region->width); - uint32_t Height = (ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY) - ? ZeImageDesc.arraylevels - : ur_cast(Region->height); - uint32_t Depth = (ZeImageDesc.type == ZE_IMAGE_TYPE_2DARRAY) - ? ZeImageDesc.arraylevels - : ur_cast(Region->depth); + uint32_t Height = ur_cast(Region->height); + uint32_t Depth = ur_cast(Region->depth); + + if (ZeImageDesc.type == ZE_IMAGE_TYPE_1D) { + OriginY = 0; + OriginZ = 0; + Height = 1; + Depth = 1; + } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY) { + // UR uses z for the array layer when describing a 1D array subregion. + // Level Zero expects the array layer in originY/height for 1D arrays. + OriginY = ur_cast(Origin->z); + OriginZ = 0; + Height = ur_cast(Region->depth); + Depth = 1; + } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D) { + OriginZ = 0; + Depth = 1; + } ZeRegion = {OriginX, OriginY, OriginZ, Width, Height, Depth}; @@ -812,6 +823,8 @@ ur_result_t bindlessImagesHandleCopyFlags( ZeStruct zeSrcImageDesc; ur2zeImageDesc(pSrcImageFormat, pSrcImageDesc, zeSrcImageDesc); + ZeStruct zeDstImageDesc; + ur2zeImageDesc(pDstImageFormat, pDstImageDesc, zeDstImageDesc); uint32_t SrcPixelSizeInBytes = getPixelSizeBytes(pSrcImageFormat); uint32_t DstPixelSizeInBytes = getPixelSizeBytes(pDstImageFormat); @@ -862,7 +875,7 @@ ur_result_t bindlessImagesHandleCopyFlags( case UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_IMAGE: { // Copy between two ze_image_handle_t's ze_image_region_t DstRegion; - UR_CALL(getZeImageRegionHelper(zeSrcImageDesc, DstPixelSizeInBytes, + UR_CALL(getZeImageRegionHelper(zeDstImageDesc, DstPixelSizeInBytes, &pCopyRegion->dstOffset, &pCopyRegion->copyExtent, DstRegion)); @@ -884,7 +897,7 @@ ur_result_t bindlessImagesHandleCopyFlags( case UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE: { // Copy from USM to ze_image_handle_t ze_image_region_t DstRegion; - UR_CALL(getZeImageRegionHelper(zeSrcImageDesc, DstPixelSizeInBytes, + UR_CALL(getZeImageRegionHelper(zeDstImageDesc, DstPixelSizeInBytes, &pCopyRegion->dstOffset, &pCopyRegion->copyExtent, DstRegion)); From f0f42ea7f098d3012f7d47c827936fc2536ad384 Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Mon, 22 Dec 2025 15:04:43 +0000 Subject: [PATCH 2/2] fix getImageRegionHelper() in l0 adapter --- .../adapters/level_zero/image_common.cpp | 60 ++++++++++--------- 1 file changed, 33 insertions(+), 27 deletions(-) diff --git a/unified-runtime/source/adapters/level_zero/image_common.cpp b/unified-runtime/source/adapters/level_zero/image_common.cpp index 6d786061ead37..dccc5673a6030 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/image_common.cpp @@ -731,33 +731,25 @@ 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); + // Runtime validation of Origin values based on image type if (ZeImageDesc.type == ZE_IMAGE_TYPE_1D) { - Region->height = 1; - Region->depth = 1; + if (Origin->y != 0 || Origin->z != 0) { + return UR_RESULT_ERROR_INVALID_VALUE; + } + } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY) { + if (Origin->y != 0) { + return UR_RESULT_ERROR_INVALID_VALUE; + } } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D) { - Region->depth = 1; + if (Origin->z != 0) { + return UR_RESULT_ERROR_INVALID_VALUE; + } } -#ifndef NDEBUG - UR_ASSERT((ZeImageDesc.type == ZE_IMAGE_TYPE_1D && Origin->y == 0 && - 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) || - (ZeImageDesc.type == ZE_IMAGE_TYPE_3D), - UR_RESULT_ERROR_INVALID_VALUE); - - UR_ASSERT(Region->width && Region->height && Region->depth, - 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->height == 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 + // Verify Region width is non-zero + if (Region->width == 0) { + return UR_RESULT_ERROR_INVALID_VALUE; + } uint32_t OriginX = ur_cast(Origin->x); uint32_t OriginY = ur_cast(Origin->y); @@ -767,9 +759,8 @@ ur_result_t getImageRegionHelper(ze_image_desc_t ZeImageDesc, uint32_t Height = ur_cast(Region->height); uint32_t Depth = ur_cast(Region->depth); + // Normalize Region dimensions based on image type if (ZeImageDesc.type == ZE_IMAGE_TYPE_1D) { - OriginY = 0; - OriginZ = 0; Height = 1; Depth = 1; } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_1DARRAY) { @@ -777,15 +768,30 @@ ur_result_t getImageRegionHelper(ze_image_desc_t ZeImageDesc, // Level Zero expects the array layer in originY/height for 1D arrays. OriginY = ur_cast(Origin->z); OriginZ = 0; - Height = ur_cast(Region->depth); + Height = ur_cast(Region->depth); // Array layer count Depth = 1; + fprintf(stderr, "[DEBUG] 1DARRAY normalization: OriginY=%u (from z), OriginZ=0, Height=%u (array layers), Depth=1\n", + OriginY, Height); } else if (ZeImageDesc.type == ZE_IMAGE_TYPE_2D) { - OriginZ = 0; Depth = 1; } ZeRegion = {OriginX, OriginY, OriginZ, Width, Height, Depth}; +#ifndef NDEBUG + // Post-normalization assertions + UR_ASSERT((ZeImageDesc.type == ZE_IMAGE_TYPE_1D && Origin->y == 0 && + 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) || + (ZeImageDesc.type == ZE_IMAGE_TYPE_3D), + UR_RESULT_ERROR_INVALID_VALUE); + + UR_ASSERT(ZeRegion.width && ZeRegion.height && ZeRegion.depth, + UR_RESULT_ERROR_INVALID_VALUE); +#endif // !NDEBUG + return UR_RESULT_SUCCESS; }