[SYCL] Fix and enable bindless image copy tests on L0

sycl/source/detail/cg.hpp

@@ -494,6 +494,7 @@ class CGCopyImage : public CG {
   ur_image_format_t MSrcImageFormat;
   ur_image_format_t MDstImageFormat;
   ur_exp_image_copy_flags_t MImageCopyFlags;
+  ur_exp_image_copy_input_types_t MImageInputTypes;
   ur_rect_offset_t MSrcOffset;
   ur_rect_offset_t MDstOffset;
   ur_rect_region_t MCopyExtent;
@@ -503,14 +504,15 @@ class CGCopyImage : public CG {
               ur_image_desc_t DstImageDesc, ur_image_format_t SrcImageFormat,
               ur_image_format_t DstImageFormat,
               ur_exp_image_copy_flags_t ImageCopyFlags,
+              ur_exp_image_copy_input_types_t ImageInputTypes,
               ur_rect_offset_t SrcOffset, ur_rect_offset_t DstOffset,
               ur_rect_region_t CopyExtent, CG::StorageInitHelper CGData,
               detail::code_location loc = {})
       : CG(CGType::CopyImage, std::move(CGData), std::move(loc)), MSrc(Src),
         MDst(Dst), MSrcImageDesc(SrcImageDesc), MDstImageDesc(DstImageDesc),
         MSrcImageFormat(SrcImageFormat), MDstImageFormat(DstImageFormat),
-        MImageCopyFlags(ImageCopyFlags), MSrcOffset(SrcOffset),
-        MDstOffset(DstOffset), MCopyExtent(CopyExtent) {}
+        MImageCopyFlags(ImageCopyFlags), MImageInputTypes(ImageInputTypes),
+        MSrcOffset(SrcOffset), MDstOffset(DstOffset), MCopyExtent(CopyExtent) {}
 
   void *getSrc() const { return MSrc; }
   void *getDst() const { return MDst; }
@@ -519,6 +521,9 @@ class CGCopyImage : public CG {
   ur_image_format_t getSrcFormat() const { return MSrcImageFormat; }
   ur_image_format_t getDstFormat() const { return MDstImageFormat; }
   ur_exp_image_copy_flags_t getCopyFlags() const { return MImageCopyFlags; }
+  ur_exp_image_copy_input_types_t getCopyInputTypes() const {
+    return MImageInputTypes;
+  }
   ur_rect_offset_t getSrcOffset() const { return MSrcOffset; }
   ur_rect_offset_t getDstOffset() const { return MDstOffset; }
   ur_rect_region_t getCopyExtent() const { return MCopyExtent; }

sycl/source/detail/handler_impl.hpp

@@ -114,6 +114,7 @@ class handler_impl {
   ur_image_format_t MSrcImageFormat = {};
   ur_image_format_t MDstImageFormat = {};
   ur_exp_image_copy_flags_t MImageCopyFlags = {};
+  ur_exp_image_copy_input_types_t MImageCopyInputTypes = {};
 
   ur_rect_offset_t MSrcOffset = {};
   ur_rect_offset_t MDestOffset = {};

sycl/source/detail/memory_manager.cpp

@@ -1570,8 +1570,10 @@ void MemoryManager::copy_image_bindless(
     queue_impl &Queue, const void *Src, void *Dst,
     const ur_image_desc_t &SrcDesc, const ur_image_desc_t &DstDesc,
     const ur_image_format_t &SrcFormat, const ur_image_format_t &DstFormat,
-    const ur_exp_image_copy_flags_t Flags, ur_rect_offset_t SrcOffset,
-    ur_rect_offset_t DstOffset, ur_rect_region_t CopyExtent,
+    const ur_exp_image_copy_flags_t Flags,
+    const ur_exp_image_copy_input_types_t InputTypes,
+    ur_rect_offset_t SrcOffset, ur_rect_offset_t DstOffset,
+    ur_rect_region_t CopyExtent,
     const std::vector<ur_event_handle_t> &DepEvents,
     ur_event_handle_t *OutEvent) {
   assert((Flags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE ||
@@ -1594,8 +1596,8 @@ void MemoryManager::copy_image_bindless(
 
   Adapter.call<UrApiKind::urBindlessImagesImageCopyExp>(
       Queue.getHandleRef(), Src, Dst, &SrcDesc, &DstDesc, &SrcFormat,
-      &DstFormat, &CopyRegion, Flags, DepEvents.size(), DepEvents.data(),
-      OutEvent);
+      &DstFormat, &CopyRegion, Flags, InputTypes, DepEvents.size(),
+      DepEvents.data(), OutEvent);
 }
 
 } // namespace detail

sycl/source/detail/memory_manager.hpp

@@ -264,8 +264,10 @@ class MemoryManager {
       queue_impl &Queue, const void *Src, void *Dst,
       const ur_image_desc_t &SrcDesc, const ur_image_desc_t &DstDesc,
       const ur_image_format_t &SrcFormat, const ur_image_format_t &DstFormat,
-      const ur_exp_image_copy_flags_t Flags, ur_rect_offset_t SrcOffset,
-      ur_rect_offset_t DstOffset, ur_rect_region_t CopyExtent,
+      const ur_exp_image_copy_flags_t Flags,
+      const ur_exp_image_copy_input_types_t InputTypes,
+      ur_rect_offset_t SrcOffset, ur_rect_offset_t DstOffset,
+      ur_rect_region_t CopyExtent,
       const std::vector<ur_event_handle_t> &DepEvents,
       ur_event_handle_t *OutEvent);
 };

sycl/source/detail/scheduler/commands.cpp

@@ -3664,8 +3664,9 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
             MemoryManager::copy_image_bindless, *MQueue, Copy->getSrc(),
             Copy->getDst(), Copy->getSrcDesc(), Copy->getDstDesc(),
             Copy->getSrcFormat(), Copy->getDstFormat(), Copy->getCopyFlags(),
-            Copy->getSrcOffset(), Copy->getDstOffset(), Copy->getCopyExtent(),
-            std::move(RawEvents), Event);
+            Copy->getCopyInputTypes(), Copy->getSrcOffset(),
+            Copy->getDstOffset(), Copy->getCopyExtent(), std::move(RawEvents),
+            Event);
         Result != UR_RESULT_SUCCESS)
       return Result;
 

sycl/source/handler.cpp

@@ -239,16 +239,16 @@ fill_image_desc(const ext::oneapi::experimental::image_descriptor &ImgDesc) {
   return UrDesc;
 }
 
-static void
-fill_copy_args(detail::handler_impl *impl,
-               const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
-               const ext::oneapi::experimental::image_descriptor &DestImgDesc,
-               ur_exp_image_copy_flags_t ImageCopyFlags, size_t SrcPitch,
-               size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
-               sycl::range<3> SrcExtent = {0, 0, 0},
-               sycl::range<3> DestOffset = {0, 0, 0},
-               sycl::range<3> DestExtent = {0, 0, 0},
-               sycl::range<3> CopyExtent = {0, 0, 0}) {
+static void fill_copy_args(
+    detail::handler_impl *impl,
+    const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
+    const ext::oneapi::experimental::image_descriptor &DestImgDesc,
+    ur_exp_image_copy_flags_t ImageCopyFlags,
+    ur_exp_image_copy_input_types_t ImageCopyInputTypes, size_t SrcPitch,
+    size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
+    sycl::range<3> SrcExtent = {0, 0, 0}, sycl::range<3> DestOffset = {0, 0, 0},
+    sycl::range<3> DestExtent = {0, 0, 0},
+    sycl::range<3> CopyExtent = {0, 0, 0}) {
   SrcImgDesc.verify();
   DestImgDesc.verify();
 
@@ -267,12 +267,13 @@ fill_copy_args(detail::handler_impl *impl,
   auto ZCopyExtentComponent = detail::fill_image_type(SrcImgDesc, UrSrcDesc);
   detail::fill_image_type(DestImgDesc, UrDestDesc);
 
-  // Copy args computed here are directly passed to UR. Various offsets and
-  // extents end up passed as ur_rect_offset_t and ur_rect_region_t. Both those
-  // structs expect their first component to be in bytes, not in pixels
-  size_t SrcPixelSize = SrcImgDesc.num_channels * get_channel_size(SrcImgDesc);
-  size_t DestPixelSize =
-      DestImgDesc.num_channels * get_channel_size(DestImgDesc);
+  // ur_rect_offset_t and ur_rect_offset_t which represent image offsets and
+  // copy extents expect that X-axis offset and region width are specified in
+  // bytes rather then in elements.
+  auto SrcPixelSize =
+      SrcImgDesc.num_channels * detail::get_channel_size(SrcImgDesc);
+  auto DestPixelSize =
+      DestImgDesc.num_channels * detail::get_channel_size(DestImgDesc);
 
   impl->MSrcOffset = {SrcOffset[0] * SrcPixelSize, SrcOffset[1], SrcOffset[2]};
   impl->MDestOffset = {DestOffset[0] * DestPixelSize, DestOffset[1],
@@ -282,6 +283,7 @@ fill_copy_args(detail::handler_impl *impl,
   impl->MSrcImageFormat = UrSrcFormat;
   impl->MDstImageFormat = UrDestFormat;
   impl->MImageCopyFlags = ImageCopyFlags;
+  impl->MImageCopyInputTypes = ImageCopyInputTypes;
 
   if (CopyExtent.size() != 0) {
     impl->MCopyExtent = {CopyExtent[0] * SrcPixelSize, CopyExtent[1],
@@ -311,6 +313,7 @@ static void
 fill_copy_args(detail::handler_impl *impl,
                const ext::oneapi::experimental::image_descriptor &Desc,
                ur_exp_image_copy_flags_t ImageCopyFlags,
+               ur_exp_image_copy_input_types_t ImageCopyInputTypes,
                sycl::range<3> SrcOffset = {0, 0, 0},
                sycl::range<3> SrcExtent = {0, 0, 0},
                sycl::range<3> DestOffset = {0, 0, 0},
@@ -320,29 +323,32 @@ fill_copy_args(detail::handler_impl *impl,
   size_t SrcPitch = SrcExtent[0] * Desc.num_channels * get_channel_size(Desc);
   size_t DestPitch = DestExtent[0] * Desc.num_channels * get_channel_size(Desc);
 
-  fill_copy_args(impl, Desc, Desc, ImageCopyFlags, SrcPitch, DestPitch,
-                 SrcOffset, SrcExtent, DestOffset, DestExtent, CopyExtent);
+  fill_copy_args(impl, Desc, Desc, ImageCopyFlags, ImageCopyInputTypes,
+                 SrcPitch, DestPitch, SrcOffset, SrcExtent, DestOffset,
+                 DestExtent, CopyExtent);
 }
 
-static void
-fill_copy_args(detail::handler_impl *impl,
-               const ext::oneapi::experimental::image_descriptor &Desc,
-               ur_exp_image_copy_flags_t ImageCopyFlags, size_t SrcPitch,
-               size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
-               sycl::range<3> SrcExtent = {0, 0, 0},
-               sycl::range<3> DestOffset = {0, 0, 0},
-               sycl::range<3> DestExtent = {0, 0, 0},
-               sycl::range<3> CopyExtent = {0, 0, 0}) {
+static void fill_copy_args(
+    detail::handler_impl *impl,
+    const ext::oneapi::experimental::image_descriptor &Desc,
+    ur_exp_image_copy_flags_t ImageCopyFlags,
+    ur_exp_image_copy_input_types_t ImageCopyInputTypes, size_t SrcPitch,
+    size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
+    sycl::range<3> SrcExtent = {0, 0, 0}, sycl::range<3> DestOffset = {0, 0, 0},
+    sycl::range<3> DestExtent = {0, 0, 0},
+    sycl::range<3> CopyExtent = {0, 0, 0}) {
 
-  fill_copy_args(impl, Desc, Desc, ImageCopyFlags, SrcPitch, DestPitch,
-                 SrcOffset, SrcExtent, DestOffset, DestExtent, CopyExtent);
+  fill_copy_args(impl, Desc, Desc, ImageCopyFlags, ImageCopyInputTypes,
+                 SrcPitch, DestPitch, SrcOffset, SrcExtent, DestOffset,
+                 DestExtent, CopyExtent);
 }
 
 static void
 fill_copy_args(detail::handler_impl *impl,
                const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
                const ext::oneapi::experimental::image_descriptor &DestImgDesc,
                ur_exp_image_copy_flags_t ImageCopyFlags,
+               ur_exp_image_copy_input_types_t ImageCopyInputTypes,
                sycl::range<3> SrcOffset = {0, 0, 0},
                sycl::range<3> SrcExtent = {0, 0, 0},
                sycl::range<3> DestOffset = {0, 0, 0},
@@ -354,9 +360,9 @@ fill_copy_args(detail::handler_impl *impl,
   size_t DestPitch =
       DestExtent[0] * DestImgDesc.num_channels * get_channel_size(DestImgDesc);
 
-  fill_copy_args(impl, SrcImgDesc, DestImgDesc, ImageCopyFlags, SrcPitch,
-                 DestPitch, SrcOffset, SrcExtent, DestOffset, DestExtent,
-                 CopyExtent);
+  fill_copy_args(impl, SrcImgDesc, DestImgDesc, ImageCopyFlags,
+                 ImageCopyInputTypes, SrcPitch, DestPitch, SrcOffset, SrcExtent,
+                 DestOffset, DestExtent, CopyExtent);
 }
 
 } // namespace detail
@@ -886,8 +892,8 @@ event handler::finalize() {
     CommandGroup.reset(new detail::CGCopyImage(
         MSrcPtr, MDstPtr, impl->MSrcImageDesc, impl->MDstImageDesc,
         impl->MSrcImageFormat, impl->MDstImageFormat, impl->MImageCopyFlags,
-        impl->MSrcOffset, impl->MDestOffset, impl->MCopyExtent,
-        std::move(impl->CGData), MCodeLoc));
+        impl->MImageCopyInputTypes, impl->MSrcOffset, impl->MDestOffset,
+        impl->MCopyExtent, std::move(impl->CGData), MCodeLoc));
     break;
   case detail::CGType::SemaphoreWait:
     CommandGroup.reset(
@@ -1313,7 +1319,8 @@ void handler::ext_oneapi_copy(
   MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
 
   detail::fill_copy_args(get_impl(), DestImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE);
+                         UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE);
 
   setType(detail::CGType::CopyImage);
 }
@@ -1331,7 +1338,8 @@ void handler::ext_oneapi_copy(
   MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
 
   detail::fill_copy_args(get_impl(), DestImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE, SrcOffset,
+                         UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE, SrcOffset,
                          SrcExtent, DestOffset, {0, 0, 0}, CopyExtent);
 
   setType(detail::CGType::CopyImage);
@@ -1348,7 +1356,8 @@ void handler::ext_oneapi_copy(
   MDstPtr = Dest;
 
   detail::fill_copy_args(get_impl(), SrcImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST);
+                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM);
 
   setType(detail::CGType::CopyImage);
 }
@@ -1367,7 +1376,8 @@ void handler::ext_oneapi_copy(
   MDstPtr = Dest;
 
   detail::fill_copy_args(get_impl(), SrcImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST, SrcOffset,
+                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM, SrcOffset,
                          {0, 0, 0}, DestOffset, DestExtent, CopyExtent);
 
   setType(detail::CGType::CopyImage);
@@ -1396,11 +1406,13 @@ 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,
-                           DeviceRowPitch);
+    detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags,
+                           UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
+                           HostRowPitch, DeviceRowPitch);
   } else if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) {
-    detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, DeviceRowPitch,
-                           HostRowPitch);
+    detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags,
+                           UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
+                           DeviceRowPitch, HostRowPitch);
   } else {
     throw sycl::exception(make_error_code(errc::invalid),
                           "Copy Error: This copy function only performs host "
@@ -1437,10 +1449,12 @@ 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,
+                           UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
                            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,
+                           UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
                            DeviceRowPitch, HostRowPitch, SrcOffset, {0, 0, 0},
                            DestOffset, HostExtent, CopyExtent);
   } else {
@@ -1465,7 +1479,8 @@ void handler::ext_oneapi_copy(
   MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
 
   detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE);
+                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_IMAGE);
 
   setType(detail::CGType::CopyImage);
 }
@@ -1485,8 +1500,10 @@ void handler::ext_oneapi_copy(
   MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
 
   detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcOffset,
-                         {0, 0, 0}, DestOffset, {0, 0, 0}, CopyExtent);
+                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_IMAGE,
+                         SrcOffset, {0, 0, 0}, DestOffset, {0, 0, 0},
+                         CopyExtent);
 
   setType(detail::CGType::CopyImage);
 }
@@ -1504,7 +1521,8 @@ 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_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM, 0,
                          DestRowPitch);
 
   setType(detail::CGType::CopyImage);
@@ -1525,7 +1543,8 @@ 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_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM, 0,
                          DestRowPitch, SrcOffset, {0, 0, 0}, DestOffset,
                          {0, 0, 0}, CopyExtent);
 
@@ -1545,8 +1564,9 @@ void handler::ext_oneapi_copy(
   MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
 
   detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcRowPitch,
-                         0);
+                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE,
+                         SrcRowPitch, 0);
 
   setType(detail::CGType::CopyImage);
 }
@@ -1566,9 +1586,10 @@ void handler::ext_oneapi_copy(
   MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
 
   detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
-                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcRowPitch,
-                         0, SrcOffset, {0, 0, 0}, DestOffset, {0, 0, 0},
-                         CopyExtent);
+                         UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
+                         UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE,
+                         SrcRowPitch, 0, SrcOffset, {0, 0, 0}, DestOffset,
+                         {0, 0, 0}, CopyExtent);
 
   setType(detail::CGType::CopyImage);
 }
@@ -1595,6 +1616,7 @@ void handler::ext_oneapi_copy(
   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,
+                           UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
                            SrcRowPitch, DestRowPitch);
   } else {
     throw sycl::exception(make_error_code(errc::invalid),
@@ -1624,6 +1646,7 @@ void handler::ext_oneapi_copy(
   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,
+                           UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
                            SrcRowPitch, DestRowPitch, SrcOffset, {0, 0, 0},
                            DestOffset, {0, 0, 0}, CopyExtent);
   } else {

sycl/test-e2e/bindless_images/copies/copy_subregion_1D.cpp

@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_oneapi_bindless_images
-// REQUIRES: cuda
 // XFAIL: hip
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/19957
 

sycl/test-e2e/bindless_images/copies/copy_subregion_2D.cpp

@@ -1,5 +1,9 @@
 // REQUIRES: aspect-ext_oneapi_bindless_images
-// REQUIRES: cuda
+// UNSUPPORTED: linux && arch-intel_gpu_bmg_g21
+// UNSUPPORTED-INTENDED: sporadic failure in CI
+//                       https://github.com/intel/llvm/issues/20006
+// XFAIL: linux && arch-intel_gpu_acm_g10
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20004
 // XFAIL: hip
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/19957