diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 6ab8a9fd03a85..5b4a4817d0da6 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -204,10 +204,14 @@ 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}; + // If this is a multi-layer array image, use the layer count; otherwise, use + // 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)) { 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..21ed10601234b 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/image_common.cpp @@ -731,34 +731,26 @@ 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) || (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(Origin->x); @@ -766,12 +758,24 @@ 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); + + // 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};