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..dccc5673a6030 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/image_common.cpp @@ -731,50 +731,67 @@ 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; + // Runtime validation of Origin values based on image type + if (ZeImageDesc.type == ZE_IMAGE_TYPE_1D) { + 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) { + if (Origin->z != 0) { + return UR_RESULT_ERROR_INVALID_VALUE; + } + } + + // 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); + uint32_t OriginZ = ur_cast(Origin->z); + + uint32_t Width = ur_cast(Region->width); + 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) { + 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); // 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) { + 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->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_ASSERT(ZeRegion.width && ZeRegion.height && ZeRegion.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->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(Origin->x); - uint32_t OriginY = ur_cast(Origin->y); - 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); - - ZeRegion = {OriginX, OriginY, OriginZ, Width, Height, Depth}; - return UR_RESULT_SUCCESS; } @@ -812,6 +829,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 +881,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 +903,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));