fix: do not use memcpy impl, match sycl docs better

Author: jharlow-intel

dpctl/_backend.pxd

@@ -511,6 +511,18 @@ cdef extern from "syclinterface/dpctl_sycl_queue_interface.h":
         size_t Count,
         const DPCTLSyclEventRef *depEvents,
         size_t depEventsCount)
+    cdef DPCTLSyclEventRef DPCTLQueue_CopyData(
+        const DPCTLSyclQueueRef Q,
+        void *Dest,
+        const void *Src,
+        size_t Count)
+    cdef DPCTLSyclEventRef DPCTLQueue_CopyDataWithEvents(
+        const DPCTLSyclQueueRef Q,
+        void *Dest,
+        const void *Src,
+        size_t Count,
+        const DPCTLSyclEventRef *depEvents,
+        size_t depEventsCount)
     cdef DPCTLSyclEventRef DPCTLQueue_Memset(
         const DPCTLSyclQueueRef Q,
         void *Dest,

dpctl/_sycl_queue.pyx

@@ -35,6 +35,8 @@ from ._backend cimport (  # noqa: E211
     DPCTLFilterSelector_Create,
     DPCTLQueue_AreEq,
     DPCTLQueue_Copy,
+    DPCTLQueue_CopyData,
+    DPCTLQueue_CopyDataWithEvents,
     DPCTLQueue_Create,
     DPCTLQueue_Delete,
     DPCTLQueue_GetBackend,
@@ -535,6 +537,80 @@ cdef DPCTLSyclEventRef _memcpy_impl(
     return ERef
 
 
+cdef DPCTLSyclEventRef _copy_impl(
+     SyclQueue q,
+     object dst,
+     object src,
+     size_t byte_count,
+     DPCTLSyclEventRef *dep_events,
+     size_t dep_events_count
+) except *:
+    cdef void *c_dst_ptr = NULL
+    cdef void *c_src_ptr = NULL
+    cdef DPCTLSyclEventRef ERef = NULL
+    cdef Py_buffer src_buf_view
+    cdef Py_buffer dst_buf_view
+    cdef bint src_is_buf = False
+    cdef bint dst_is_buf = False
+    cdef int ret_code = 0
+
+    if isinstance(src, _Memory):
+        c_src_ptr = <void*>(<_Memory>src).get_data_ptr()
+    elif _is_buffer(src):
+        ret_code = PyObject_GetBuffer(
+            src, &src_buf_view, PyBUF_SIMPLE | PyBUF_ANY_CONTIGUOUS
+        )
+        if ret_code != 0:  # pragma: no cover
+            raise RuntimeError("Could not access buffer")
+        c_src_ptr = src_buf_view.buf
+        src_is_buf = True
+    else:
+        raise TypeError(
+             "Parameter `src` should have either type "
+             "`dpctl.memory._Memory` or a type that "
+             "supports Python buffer protocol"
+        )
+
+    if isinstance(dst, _Memory):
+        c_dst_ptr = <void*>(<_Memory>dst).get_data_ptr()
+    elif _is_buffer(dst):
+        ret_code = PyObject_GetBuffer(
+            dst, &dst_buf_view,
+            PyBUF_SIMPLE | PyBUF_ANY_CONTIGUOUS | PyBUF_WRITABLE
+        )
+        if ret_code != 0:  # pragma: no cover
+            if src_is_buf:
+                PyBuffer_Release(&src_buf_view)
+            raise RuntimeError("Could not access buffer")
+        c_dst_ptr = dst_buf_view.buf
+        dst_is_buf = True
+    else:
+        raise TypeError(
+             "Parameter `dst` should have either type "
+             "`dpctl.memory._Memory` or a type that "
+             "supports Python buffer protocol"
+        )
+
+    if dep_events_count == 0 or dep_events is NULL:
+        ERef = DPCTLQueue_CopyData(q._queue_ref, c_dst_ptr, c_src_ptr, byte_count)
+    else:
+        ERef = DPCTLQueue_CopyDataWithEvents(
+            q._queue_ref,
+            c_dst_ptr,
+            c_src_ptr,
+            byte_count,
+            dep_events,
+            dep_events_count
+        )
+
+    if src_is_buf:
+        PyBuffer_Release(&src_buf_view)
+    if dst_is_buf:
+        PyBuffer_Release(&dst_buf_view)
+
+    return ERef
+
+
 cdef class _SyclQueue:
     """ Barebone data owner class used by SyclQueue.
     """
@@ -1381,7 +1457,7 @@ cdef class SyclQueue(_SyclQueue):
         """Copy memory from `src` to `dst`"""
         cdef DPCTLSyclEventRef ERef = NULL
 
-        ERef = _memcpy_impl(<SyclQueue>self, dest, src, count, NULL, 0)
+        ERef = _copy_impl(<SyclQueue>self, dest, src, count, NULL, 0)
         if (ERef is NULL):
             raise RuntimeError(
                 "SyclQueue.memcpy operation encountered an error"
@@ -1399,7 +1475,7 @@ cdef class SyclQueue(_SyclQueue):
         cdef size_t nDE = 0
 
         if dEvents is None:
-            ERef = _memcpy_impl(<SyclQueue>self, dest, src, count, NULL, 0)
+            ERef = _copy_impl(<SyclQueue>self, dest, src, count, NULL, 0)
         else:
             nDE = len(dEvents)
             depEvents = (
@@ -1416,7 +1492,7 @@ cdef class SyclQueue(_SyclQueue):
                         raise TypeError(
                             "A sequence of dpctl.SyclEvent is expected"
                         )
-            ERef = _memcpy_impl(self, dest, src, count, depEvents, nDE)
+            ERef = _copy_impl(self, dest, src, count, depEvents, nDE)
             free(depEvents)
 
         if (ERef is NULL):
@@ -1429,6 +1505,9 @@ cdef class SyclQueue(_SyclQueue):
     cpdef copy(self, dest, src, size_t count):
         """Copy ``count`` bytes from ``src`` to ``dest`` and wait.
 
+        Internally, this dispatches ``sycl::queue::copy`` instantiated for
+        byte-sized elements.
+
         This is a synchronizing variant corresponding to
         :meth:`dpctl.SyclQueue.copy_async`.
         """
@@ -1448,6 +1527,9 @@ cdef class SyclQueue(_SyclQueue):
     ):
         """Copy ``count`` bytes from ``src`` to ``dest`` asynchronously.
 
+        Internally, this dispatches ``sycl::queue::copy`` instantiated for
+        byte-sized elements.
+
         Args:
             dest:
                 Destination USM object or Python object supporting

dpctl/tests/test_sycl_queue_memcpy.py

@@ -155,6 +155,40 @@ def test_memcpy_async():
     assert dst_buf2 == src_buf
 
 
+def test_copy_copy_host_to_host():
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Default constructor for SyclQueue failed")
+
+    src_buf = b"abcdefghijklmnopqrstuvwxyz"
+    dst_buf = bytearray(len(src_buf))
+
+    q.copy(dst_buf, src_buf, len(src_buf))
+
+    assert dst_buf == src_buf
+
+
+def test_copy_async():
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Default constructor for SyclQueue failed")
+
+    src_buf = b"abcdefghijklmnopqrstuvwxyz"
+    n = len(src_buf)
+    dst_buf = bytearray(n)
+    dst_buf2 = bytearray(n)
+
+    e = q.copy_async(dst_buf, src_buf, n)
+    e2 = q.copy_async(dst_buf2, src_buf, n, [e])
+
+    e.wait()
+    e2.wait()
+    assert dst_buf == src_buf
+    assert dst_buf2 == src_buf
+
+
 def test_memcpy_type_error():
     try:
         q = dpctl.SyclQueue()

libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h

@@ -292,7 +292,9 @@ void DPCTLQueue_Wait(__dpctl_keep const DPCTLSyclQueueRef QRef);
  * @param    QRef           An opaque pointer to the ``sycl::queue``.
  * @param    Dest           An USM pointer to the destination memory.
  * @param    Src            An USM pointer to the source memory.
- * @param    Count          A number of bytes to copy.
+ * @param    Count          A number of bytes to copy. The wrapper binds
+ *                          ``sycl::queue::copy`` with ``T=uint8_t`` so the
+ *                          SYCL element count matches a byte count.
  * @return   An opaque pointer to the ``sycl::event`` returned by the
  *           ``sycl::queue::memcpy`` function.
  * @ingroup QueueInterface
@@ -310,7 +312,9 @@ DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
  * @param    QRef           An opaque pointer to the ``sycl::queue``.
  * @param    Dest           An USM pointer to the destination memory.
  * @param    Src            An USM pointer to the source memory.
- * @param    Count          A number of bytes to copy.
+ * @param    Count          A number of bytes to copy. The wrapper binds
+ *                          ``sycl::queue::copy`` with ``T=uint8_t`` so the
+ *                          SYCL element count matches a byte count.
  * @param    DepEvents      A pointer to array of DPCTLSyclEventRef opaque
  *                          pointers to dependent events.
  * @param    DepEventsCount A number of dependent events.
@@ -327,6 +331,47 @@ DPCTLQueue_MemcpyWithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef,
                             __dpctl_keep const DPCTLSyclEventRef *DepEvents,
                             size_t DepEventsCount);
 
+/*!
+ * @brief C-API wrapper for ``sycl::queue::copy``.
+ *
+ * @param    QRef           An opaque pointer to the ``sycl::queue``.
+ * @param    Dest           A destination pointer.
+ * @param    Src            A source pointer.
+ * @param    Count          A number of bytes to copy.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::copy`` function.
+ * @ingroup QueueInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_CopyData(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                    void *Dest,
+                    const void *Src,
+                    size_t Count);
+
+/*!
+ * @brief C-API wrapper for ``sycl::queue::copy``.
+ *
+ * @param    QRef           An opaque pointer to the ``sycl::queue``.
+ * @param    Dest           A destination pointer.
+ * @param    Src            A source pointer.
+ * @param    Count          A number of bytes to copy.
+ * @param    DepEvents      A pointer to array of DPCTLSyclEventRef opaque
+ *                          pointers to dependent events.
+ * @param    DepEventsCount A number of dependent events.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::copy`` function.
+ * @ingroup QueueInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_CopyDataWithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                              void *Dest,
+                              const void *Src,
+                              size_t Count,
+                              __dpctl_keep const DPCTLSyclEventRef *DepEvents,
+                              size_t DepEventsCount);
+
 /*!
  * @brief C-API wrapper for ``sycl::queue::prefetch``.
  *

libsyclinterface/source/dpctl_sycl_queue_interface.cpp

@@ -41,6 +41,7 @@
 #include <stdexcept>
 #include <sycl/sycl.hpp> /* SYCL headers   */
 #include <utility>
+#include <vector>
 
 #if defined(SYCL_EXT_ONEAPI_WORK_GROUP_MEMORY) ||                              \
     defined(SYCL_EXT_ONEAPI_RAW_KERNEL_ARG)
@@ -693,6 +694,70 @@ DPCTLQueue_MemcpyWithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef,
     return wrap<event>(new event(ev));
 }
 
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_CopyData(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                    void *Dest,
+                    const void *Src,
+                    size_t Count)
+{
+    auto Q = unwrap<queue>(QRef);
+    if (Q) {
+        sycl::event ev;
+        try {
+            // Bind queue::copy with uint8_t so Count is interpreted as bytes.
+            ev = Q->copy(static_cast<const std::uint8_t *>(Src),
+                         static_cast<std::uint8_t *>(Dest), Count);
+        } catch (std::exception const &e) {
+            error_handler(e, __FILE__, __func__, __LINE__);
+            return nullptr;
+        }
+        return wrap<event>(new event(std::move(ev)));
+    }
+    else {
+        error_handler("QRef passed to copy was NULL.", __FILE__, __func__,
+                      __LINE__);
+        return nullptr;
+    }
+}
+
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_CopyDataWithEvents(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                              void *Dest,
+                              const void *Src,
+                              size_t Count,
+                              const DPCTLSyclEventRef *DepEvents,
+                              size_t DepEventsCount)
+{
+    auto Q = unwrap<queue>(QRef);
+    if (Q) {
+        try {
+            std::vector<event> dep_events;
+            if (DepEvents) {
+                dep_events.reserve(DepEventsCount);
+                for (size_t i = 0; i < DepEventsCount; ++i) {
+                    event *ei = unwrap<event>(DepEvents[i]);
+                    if (ei)
+                        dep_events.push_back(*ei);
+                }
+            }
+
+            // Bind queue::copy with uint8_t so Count is interpreted as bytes.
+            auto ev = Q->copy(static_cast<const std::uint8_t *>(Src),
+                              static_cast<std::uint8_t *>(Dest), Count,
+                              dep_events);
+            return wrap<event>(new event(std::move(ev)));
+        } catch (const std::exception &ex) {
+            error_handler(ex, __FILE__, __func__, __LINE__);
+            return nullptr;
+        }
+    }
+    else {
+        error_handler("QRef passed to copy was NULL.", __FILE__, __func__,
+                      __LINE__);
+        return nullptr;
+    }
+}
+
 __dpctl_give DPCTLSyclEventRef
 DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
                     const void *Ptr,

libsyclinterface/tests/test_sycl_queue_interface.cpp

@@ -371,6 +371,13 @@ TEST(TestDPCTLSyclQueueInterface, CheckMemOpsZeroQRef)
         ERef = DPCTLQueue_MemcpyWithEvents(QRef, p1, p2, n_bytes, NULL, 0));
     ASSERT_FALSE(bool(ERef));
 
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_CopyData(QRef, p1, p2, n_bytes));
+    ASSERT_FALSE(bool(ERef));
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_CopyDataWithEvents(
+                                QRef, p1, p2, n_bytes, NULL, 0));
+    ASSERT_FALSE(bool(ERef));
+
     ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Prefetch(QRef, p1, n_bytes));
     ASSERT_FALSE(bool(ERef));
 
@@ -429,6 +436,13 @@ TEST_P(TestDPCTLQueueMemberFunctions, CheckMemOpsNullPtr)
         ERef = DPCTLQueue_MemcpyWithEvents(QRef, p1, p2, n_bytes, NULL, 0));
     ASSERT_FALSE(bool(ERef));
 
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_CopyData(QRef, p1, p2, n_bytes));
+    ASSERT_FALSE(bool(ERef));
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_CopyDataWithEvents(
+                                QRef, p1, p2, n_bytes, NULL, 0));
+    ASSERT_FALSE(bool(ERef));
+
     ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Prefetch(QRef, p1, n_bytes));
     if (ERef) {
         ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(ERef));