[SYCL][Bindless] Add image_mem_handle to image_mem_handle devices copies.

Author: cppchedy

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -749,6 +749,12 @@ public:
     size_t DeviceRowPitch,
     sycl::range<3> HostExtent,
     sycl::range<3> CopyExtent);
+
+  // Simple device to device copy
+  void ext_oneapi_copy(
+      ext::oneapi::experimental::image_mem_handle Src,
+      ext::oneapi::experimental::image_mem_handle Dest,
+      const ext::oneapi::experimental::image_descriptor &ImageDesc);
 };
 
 class queue {
@@ -874,14 +880,32 @@ public:
       size_t DeviceRowPitch,
       sycl::range<3> HostExtent,
       sycl::range<3> CopyExtent);
+  
+  // Simple device to device copy
+  event ext_oneapi_copy(
+      ext::oneapi::experimental::image_mem_handle Src,
+      ext::oneapi::experimental::image_mem_handle Dest,
+      const ext::oneapi::experimental::image_descriptor &ImageDesc);
+  event ext_oneapi_copy(
+      ext::oneapi::experimental::image_mem_handle Src,
+      ext::oneapi::experimental::image_mem_handle Dest,
+      const ext::oneapi::experimental::image_descriptor &ImageDesc,
+      event DepEvent);
+  event ext_oneapi_copy(
+      ext::oneapi::experimental::image_mem_handle Src,
+      ext::oneapi::experimental::image_mem_handle Dest,
+      const ext::oneapi::experimental::image_descriptor &ImageDesc,
+      const std::vector<event> &DepEvents);
 };
 }
 ```
 
 To enable the copying of images an `ext_oneapi_copy` function is proposed as a 
 method of the queue and handler. It can be used to copy image memory, whether 
 allocated through USM or using an `image_mem_handle`, from host to 
-device, or device to host. For the `ext_oneapi_copy` variants that do not take 
+device, or device to host. Device to device copies are currently supported only 
+through `image_mem_handle` allocations. 
+For the `ext_oneapi_copy` variants that do not take 
 offsets and extents, the image descriptor passed to the `ext_oneapi_copy` API 
 is used to determine the pixel size, dimensions, and extent in memory of the 
 image to copy. If performing sub-region copy, the size of the memory region is 
@@ -2060,4 +2084,6 @@ These features still need to be handled:
                    wording around what types are allowed to be read or written.
                  - Allow `read_image` and `read_mipmap` to return a 
                    user-defined type.
+|5.1|2024-01-17| - Added overload for `ext_oneapi_copy` enabling device to device
+                   copies using `image_mem_handle`.
 |======================

sycl/include/sycl/handler.hpp

@@ -3227,6 +3227,18 @@ class __SYCL_EXPORT handler {
       const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
       size_t DeviceRowPitch);
 
+  /// Copies data from device to device memory, where \p Src and \p Dest
+  /// are opaque image memory handles.
+  /// An exception is thrown if either \p Src or \p Dest is incomplete
+  ///
+  /// \param Src is an opaque image memory handle to the source memory.
+  /// \param Dest is an opaque image memory handle to the destination memory.
+  /// \param ImageDesc is the source image descriptor
+  void
+  ext_oneapi_copy(ext::oneapi::experimental::image_mem_handle Src,
+                  ext::oneapi::experimental::image_mem_handle Dest,
+                  const ext::oneapi::experimental::image_descriptor &ImageDesc);
+
   /// Copies data from one memory region to another, where \p Src and \p Dest
   /// are USM pointers. Allows for a sub-region copy, where \p SrcOffset ,
   /// \p DestOffset , and \p Extent are used to determine the sub-region.

sycl/include/sycl/queue.hpp

@@ -1917,6 +1917,74 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
         CodeLoc);
   }
 
+  /// Copies data from device to device memory, where \p Src and \p Dest
+  /// are opaque image memory handles.
+  /// An exception is thrown if either \p Src or \p Dest is incomplete
+  ///
+  /// \param Src is an opaque image memory handle to the source memory.
+  /// \param Dest is an opaque image memory handle to the destination memory.
+  /// \param ImageDesc is the source image descriptor
+  /// \param DepEvent is an events that specifies the kernel dependency.
+  /// \return an event representing the copy operation.
+  event ext_oneapi_copy(
+      ext::oneapi::experimental::image_mem_handle Src,
+      ext::oneapi::experimental::image_mem_handle Dest,
+      const ext::oneapi::experimental::image_descriptor &ImageDesc,
+      event DepEvent,
+      const detail::code_location &CodeLoc = detail::code_location::current()) {
+    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+    return submit(
+        [&](handler &CGH) {
+          CGH.depends_on(DepEvent);
+          CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
+        },
+        CodeLoc);
+  }
+
+  /// Copies data from device to device memory, where \p Src and \p Dest
+  /// are opaque image memory handles.
+  /// An exception is thrown if either \p Src or \p Dest is incomplete
+  ///
+  /// \param Src is an opaque image memory handle to the source memory.
+  /// \param Dest is an opaque image memory handle to the destination memory.
+  /// \param ImageDesc is the source image descriptor
+  /// \param DepEvents is a vector of events that specifies the kernel
+  /// dependencies.
+  /// \return an event representing the copy operation.
+  event ext_oneapi_copy(
+      ext::oneapi::experimental::image_mem_handle Src,
+      ext::oneapi::experimental::image_mem_handle Dest,
+      const ext::oneapi::experimental::image_descriptor &ImageDesc,
+      const std::vector<event> &DepEvents,
+      const detail::code_location &CodeLoc = detail::code_location::current()) {
+    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+    return submit(
+        [&](handler &CGH) {
+          CGH.depends_on(DepEvents);
+          CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
+        },
+        CodeLoc);
+  }
+
+  /// Copies data from device to device memory, where \p Src and \p Dest
+  /// are opaque image memory handles.
+  /// An exception is thrown if either \p Src or \p Dest is incomplete
+  ///
+  /// \param Src is an opaque image memory handle to the source memory.
+  /// \param Dest is an opaque image memory handle to the destination memory.
+  /// \param ImageDesc is the source image descriptor
+  /// \return an event representing the copy operation.
+  event ext_oneapi_copy(
+      ext::oneapi::experimental::image_mem_handle Src,
+      ext::oneapi::experimental::image_mem_handle Dest,
+      const ext::oneapi::experimental::image_descriptor &ImageDesc,
+      const detail::code_location &CodeLoc = detail::code_location::current()) {
+    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+    return submit(
+        [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, ImageDesc); },
+        CodeLoc);
+  }
+
   /// Copies data from one memory region to another, where \p Src and \p Dest
   /// are USM pointers. Allows for a sub-region copy, where \p SrcOffset ,
   /// \p DestOffset , and \p Extent are used to determine the sub-region.

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -56,14 +56,9 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 79c28d0f0713f58358d5080653d95803fd131749
-  # Merge: 25e0b603 45d76b78
-  # Author: aarongreig <[email protected]>
-  # Date: Fri Jan 12 16:14:44 2024 +0000
-  #     Merge pull request #1186 from hdelan/device-global-hip
-  #     [HIP] Add support for global variable read write
-  set(UNIFIED_RUNTIME_TAG 79c28d0f0713f58358d5080653d95803fd131749)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/cppchedy/unified-runtime.git")
+  set(UNIFIED_RUNTIME_TAG 0942022ba947f1832056ffa9e317dc1384c382e0)
+
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")

sycl/source/detail/memory_manager.cpp

@@ -1755,7 +1755,9 @@ void MemoryManager::copy_image_bindless(
   assert((Flags == (sycl::detail::pi::PiImageCopyFlags)
                        ext::oneapi::experimental::image_copy_flags::HtoD ||
           Flags == (sycl::detail::pi::PiImageCopyFlags)
-                       ext::oneapi::experimental::image_copy_flags::DtoH) &&
+                       ext::oneapi::experimental::image_copy_flags::DtoH ||
+          Flags == (sycl::detail::pi::PiImageCopyFlags)
+                       ext::oneapi::experimental::image_copy_flags::DtoD) &&
          "Invalid flags passed to copy_image_bindless.");
   if (!Dst || !Src)
     throw sycl::exception(

sycl/source/handler.cpp

@@ -1058,6 +1058,42 @@ void handler::ext_oneapi_copy(
   setType(detail::CG::CopyImage);
 }
 
+void handler::ext_oneapi_copy(
+    ext::oneapi::experimental::image_mem_handle Src,
+    ext::oneapi::experimental::image_mem_handle Dest,
+    const ext::oneapi::experimental::image_descriptor &ImageDesc) {
+  throwIfGraphAssociated<
+      ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+          sycl_ext_oneapi_bindless_images>();
+  MSrcPtr = Src.raw_handle;
+  MDstPtr = Dest.raw_handle;
+
+  sycl::detail::pi::PiMemImageDesc PiDesc = {};
+  PiDesc.image_width = ImageDesc.width;
+  PiDesc.image_height = ImageDesc.height;
+  PiDesc.image_depth = ImageDesc.depth;
+  PiDesc.image_type =
+      ImageDesc.depth > 0
+          ? PI_MEM_TYPE_IMAGE3D
+          : (ImageDesc.height > 0 ? PI_MEM_TYPE_IMAGE2D : PI_MEM_TYPE_IMAGE1D);
+
+  sycl::detail::pi::PiMemImageFormat PiFormat;
+  PiFormat.image_channel_data_type =
+      sycl::_V1::detail::convertChannelType(ImageDesc.channel_type);
+  PiFormat.image_channel_order =
+      sycl::_V1::detail::convertChannelOrder(ImageDesc.channel_order);
+
+  MImpl->MSrcOffset = {0, 0, 0};
+  MImpl->MDestOffset = {0, 0, 0};
+  MImpl->MCopyExtent = {ImageDesc.width, ImageDesc.height, ImageDesc.depth};
+  MImpl->MHostExtent = {ImageDesc.width, ImageDesc.height, ImageDesc.depth};
+  MImpl->MImageDesc = PiDesc;
+  MImpl->MImageFormat = PiFormat;
+  MImpl->MImageCopyFlags =
+      sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_DEVICE_TO_DEVICE;
+  setType(detail::CG::CopyImage);
+}
+
 void handler::ext_oneapi_copy(
     ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset,
     const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,

sycl/test-e2e/bindless_images/device_to_device_copy.cpp

@@ -0,0 +1,99 @@
+// REQUIRES: linux
+// REQUIRES: cuda
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+// Uncomment to print additional test information
+#define VERBOSE_PRINT
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+void copy_image_mem_handle_to_image_mem_handle(
+    syclexp::image_descriptor &desc, const std::vector<float> &testData,
+    sycl::device dev, sycl::queue q, std::vector<float> &out) {
+  syclexp::image_mem imgMemSrc(desc, dev, q.get_context());
+  syclexp::image_mem imgMemDst(desc, dev, q.get_context());
+
+  q.ext_oneapi_copy((void *)testData.data(), imgMemSrc.get_handle(), desc);
+  q.wait_and_throw();
+
+  q.ext_oneapi_copy(imgMemSrc.get_handle(), imgMemDst.get_handle(), desc);
+  q.wait_and_throw();
+
+  q.ext_oneapi_copy(imgMemDst.get_handle(), (void *)out.data(), desc);
+  q.wait_and_throw();
+}
+
+bool check_test(const std::vector<float> &out,
+                const std::vector<float> &expected) {
+  assert(out.size() == expected.size());
+  bool validated = true;
+  for (int i = 0; i < out.size(); i++) {
+    bool mismatch = false;
+    if (out[i] != expected[i]) {
+      mismatch = true;
+      validated = false;
+    }
+
+    if (mismatch) {
+#ifdef VERBOSE_PRINT
+      std::cout << "Result mismatch! Expected: " << expected[i]
+                << ", Actual: " << out[i] << std::endl;
+#else
+      break;
+#endif
+    }
+  }
+  return validated;
+}
+
+template <sycl::image_channel_order channelOrder,
+          sycl::image_channel_type channelType, int dim>
+bool run_copy_test_with(sycl::device &dev, sycl::queue &q,
+                        sycl::range<dim> dims) {
+  std::vector<float> dataSequence(dims.size());
+  std::vector<float> out(dims.size());
+
+  std::vector<float> expected(dims.size());
+
+  std::iota(dataSequence.begin(), dataSequence.end(), 0);
+  std::iota(expected.begin(), expected.end(), 0);
+
+  syclexp::image_descriptor desc(dims, channelOrder, channelType);
+
+  copy_image_mem_handle_to_image_mem_handle(desc, dataSequence, dev, q, out);
+
+  return check_test(out, expected);
+}
+
+int main() {
+
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  bool validated =
+      run_copy_test_with<sycl::image_channel_order::r,
+                         sycl::image_channel_type::fp32, 1>(dev, q, {4});
+
+  validated &=
+      run_copy_test_with<sycl::image_channel_order::r,
+                         sycl::image_channel_type::fp32, 2>(dev, q, {4, 4});
+
+  validated &=
+      run_copy_test_with<sycl::image_channel_order::r,
+                         sycl::image_channel_type::fp32, 3>(dev, q, {4, 4, 4});
+
+  if (!validated) {
+    std::cout << "Tests failed";
+    return 1;
+  }
+
+  std::cout << "Tests passed";
+
+  return 0;
+}