diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index f9d1769e573e7..4a033eab16681 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -1661,10 +1661,12 @@ void handler::ext_oneapi_copy( Desc.width * Desc.num_channels * detail::get_channel_size(Desc); if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE) { - detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, HostRowPitch, + detail::fill_copy_args(get_impl(), Desc, + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST, HostRowPitch, DeviceRowPitch); } else if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) { - detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, DeviceRowPitch, + detail::fill_copy_args(get_impl(), Desc, + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST, DeviceRowPitch, HostRowPitch); } else { throw sycl::exception(make_error_code(errc::invalid), @@ -1701,13 +1703,15 @@ void handler::ext_oneapi_copy( // Fill the host extent based on the type of copy. if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE) { - detail::fill_copy_args(get_impl(), DeviceImgDesc, ImageCopyFlags, - HostRowPitch, DeviceRowPitch, SrcOffset, HostExtent, - DestOffset, {0, 0, 0}, CopyExtent); + detail::fill_copy_args(get_impl(), DeviceImgDesc, + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST, HostRowPitch, + DeviceRowPitch, SrcOffset, HostExtent, DestOffset, + {0, 0, 0}, CopyExtent); } else if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) { - detail::fill_copy_args(get_impl(), DeviceImgDesc, ImageCopyFlags, - DeviceRowPitch, HostRowPitch, SrcOffset, {0, 0, 0}, - DestOffset, HostExtent, CopyExtent); + detail::fill_copy_args(get_impl(), DeviceImgDesc, + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST, DeviceRowPitch, + HostRowPitch, SrcOffset, {0, 0, 0}, DestOffset, + HostExtent, CopyExtent); } else { throw sycl::exception(make_error_code(errc::invalid), "Copy Error: This copy function only performs host " @@ -1769,7 +1773,7 @@ void handler::ext_oneapi_copy( MDstPtr = Dest; detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, - UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, 0, + UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST, 0, DestRowPitch); setType(detail::CGType::CopyImage); @@ -1790,7 +1794,7 @@ void handler::ext_oneapi_copy( MDstPtr = Dest; detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, - UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, 0, + UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST, 0, DestRowPitch, SrcOffset, {0, 0, 0}, DestOffset, {0, 0, 0}, CopyExtent); @@ -1810,7 +1814,7 @@ void handler::ext_oneapi_copy( MDstPtr = reinterpret_cast(Dest.raw_handle); detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, - UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcRowPitch, + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE, SrcRowPitch, 0); setType(detail::CGType::CopyImage); @@ -1831,7 +1835,7 @@ void handler::ext_oneapi_copy( MDstPtr = reinterpret_cast(Dest.raw_handle); detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, - UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcRowPitch, + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE, SrcRowPitch, 0, SrcOffset, {0, 0, 0}, DestOffset, {0, 0, 0}, CopyExtent); @@ -1851,21 +1855,10 @@ void handler::ext_oneapi_copy( MSrcPtr = const_cast(Src); MDstPtr = Dest; - ur_exp_image_copy_flags_t ImageCopyFlags = detail::getUrImageCopyFlags( - get_pointer_type(Src, - createSyclObjFromImpl(impl->get_context())), - get_pointer_type(Dest, - createSyclObjFromImpl(impl->get_context()))); - - if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE || - ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST) { - detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, ImageCopyFlags, - SrcRowPitch, DestRowPitch); - } else { - throw sycl::exception(make_error_code(errc::invalid), - "Copy Error: This copy function only performs device " - "to device or host to host copies!"); - } + ur_exp_image_copy_flags_t ImageCopyFlags = + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST; + detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, ImageCopyFlags, + SrcRowPitch, DestRowPitch); setType(detail::CGType::CopyImage); } @@ -1880,22 +1873,11 @@ void handler::ext_oneapi_copy( MSrcPtr = const_cast(Src); MDstPtr = Dest; - ur_exp_image_copy_flags_t ImageCopyFlags = detail::getUrImageCopyFlags( - get_pointer_type(Src, - createSyclObjFromImpl(impl->get_context())), - get_pointer_type(Dest, - createSyclObjFromImpl(impl->get_context()))); - - if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE || - ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST) { - detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, ImageCopyFlags, - SrcRowPitch, DestRowPitch, SrcOffset, {0, 0, 0}, - DestOffset, {0, 0, 0}, CopyExtent); - } else { - throw sycl::exception(make_error_code(errc::invalid), - "Copy Error: This copy function only performs device " - "to device or host to host copies!"); - } + ur_exp_image_copy_flags_t ImageCopyFlags = + UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST; + detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, ImageCopyFlags, + SrcRowPitch, DestRowPitch, SrcOffset, {0, 0, 0}, + DestOffset, {0, 0, 0}, CopyExtent); setType(detail::CGType::CopyImage); } diff --git a/sycl/test-e2e/bindless_images/copies/copy_subregion_1D.cpp b/sycl/test-e2e/bindless_images/copies/copy_subregion_1D.cpp index ac93ae5db9e67..ff9430ee31fc0 100644 --- a/sycl/test-e2e/bindless_images/copies/copy_subregion_1D.cpp +++ b/sycl/test-e2e/bindless_images/copies/copy_subregion_1D.cpp @@ -1,5 +1,4 @@ // REQUIRES: aspect-ext_oneapi_bindless_images -// REQUIRES: cuda // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/bindless_images/copies/copy_subregion_2D.cpp b/sycl/test-e2e/bindless_images/copies/copy_subregion_2D.cpp index 6d61a2a2a259a..c3325fbef33b0 100644 --- a/sycl/test-e2e/bindless_images/copies/copy_subregion_2D.cpp +++ b/sycl/test-e2e/bindless_images/copies/copy_subregion_2D.cpp @@ -1,5 +1,4 @@ // REQUIRES: aspect-ext_oneapi_bindless_images -// REQUIRES: cuda // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/bindless_images/copies/device_to_device_copy.cpp b/sycl/test-e2e/bindless_images/copies/device_to_device_copy.cpp index 9777e2a591837..2b2028c5a7e06 100644 --- a/sycl/test-e2e/bindless_images/copies/device_to_device_copy.cpp +++ b/sycl/test-e2e/bindless_images/copies/device_to_device_copy.cpp @@ -1,5 +1,4 @@ // REQUIRES: aspect-ext_oneapi_bindless_images -// REQUIRES: cuda // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/bindless_images/copies/device_to_device_pitched.cpp b/sycl/test-e2e/bindless_images/copies/device_to_device_pitched.cpp index c8a644296eba2..1842cd30735b0 100644 --- a/sycl/test-e2e/bindless_images/copies/device_to_device_pitched.cpp +++ b/sycl/test-e2e/bindless_images/copies/device_to_device_pitched.cpp @@ -1,6 +1,7 @@ // REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_bindless_images_2d_usm -// REQUIRES: cuda +// XFAIL: level_zero +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/17663 // // UNSUPPORTED: cuda // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17231 diff --git a/sycl/test-e2e/bindless_images/copies/host_to_host_pitched.cpp b/sycl/test-e2e/bindless_images/copies/host_to_host_pitched.cpp index 37e58a3d75bef..41405b891b484 100644 --- a/sycl/test-e2e/bindless_images/copies/host_to_host_pitched.cpp +++ b/sycl/test-e2e/bindless_images/copies/host_to_host_pitched.cpp @@ -1,6 +1,7 @@ // REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_bindless_images_2d_usm -// REQUIRES: cuda +// XFAIL: level_zero +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/17663 // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/bindless_images/sampling_2D_USM_host.cpp b/sycl/test-e2e/bindless_images/sampling_2D_USM_host.cpp index bca3d2c1c0ddd..4e369dfef27b8 100644 --- a/sycl/test-e2e/bindless_images/sampling_2D_USM_host.cpp +++ b/sycl/test-e2e/bindless_images/sampling_2D_USM_host.cpp @@ -1,5 +1,6 @@ -// REQUIRES: cuda // REQUIRES: aspect-ext_oneapi_bindless_images_2d_usm +// XFAIL: level_zero +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/17663 // RUN: %{build} -o %t.out // RUN: %{run-unfiltered-devices} %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 9b76788b6a65e..d4601beee7f90 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/image_common.cpp @@ -882,6 +882,30 @@ ur_result_t bindlessImagesHandleCopyFlags( return UR_RESULT_SUCCESS; }; + case UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST: { + // Copy between pitched USM regions + uint32_t DstRowPitch = pDstImageDesc->rowPitch; + uint32_t SrcRowPitch = pSrcImageDesc->rowPitch; + ze_copy_region_t ZeDstRegion = {(uint32_t)pCopyRegion->dstOffset.x, + (uint32_t)pCopyRegion->dstOffset.y, + (uint32_t)pCopyRegion->dstOffset.z, + DstRowPitch, + (uint32_t)pCopyRegion->copyExtent.height, + (uint32_t)pCopyRegion->copyExtent.depth}; + ze_copy_region_t ZeSrcRegion = {(uint32_t)pCopyRegion->srcOffset.x, + (uint32_t)pCopyRegion->srcOffset.y, + (uint32_t)pCopyRegion->srcOffset.z, + SrcRowPitch, + (uint32_t)pCopyRegion->copyExtent.height, + (uint32_t)pCopyRegion->copyExtent.depth}; + uint32_t DstSlicePitch = DstRowPitch * pDstImageDesc->height; + uint32_t SrcSlicePitch = SrcRowPitch * pSrcImageDesc->height; + ZE2UR_CALL(zeCommandListAppendMemoryCopyRegion, + (ZeCommandList, pDst, &ZeDstRegion, DstRowPitch, DstSlicePitch, + pSrc, &ZeSrcRegion, SrcRowPitch, SrcSlicePitch, zeSignalEvent, + numWaitEvents, phWaitEvents)); + return UR_RESULT_SUCCESS; + }; default: UR_LOG(ERR, "ur_queue_immediate_in_order_t::bindlessImagesImageCopyExp: " "unexpected imageCopyFlags");