DPCTLQueue_Memcpy, _Prefetch, _Memadvise become asynchronous

dpctl-capi/include/dpctl_sycl_queue_interface.h

@@ -266,7 +266,7 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef,
                          size_t NDepEvents);
 
 /*!
- * @brief Calls the ``sycl::queue.submit`` function to do a blocking wait on
+ * @brief Calls the ``sycl::queue::submit`` function to do a blocking wait on
  * all enqueued tasks in the queue.
  *
  * @param    QRef           Opaque pointer to a ``sycl::queue``.
@@ -276,52 +276,55 @@ DPCTL_API
 void DPCTLQueue_Wait(__dpctl_keep const DPCTLSyclQueueRef QRef);
 
 /*!
- * @brief C-API wrapper for ``sycl::queue::memcpy``, the function waits on an
- * event till the memcpy operation completes.
+ * @brief C-API wrapper for ``sycl::queue::memcpy``.
  *
  * @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.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::memcpy`` function.
  * @ingroup QueueInterface
  */
 DPCTL_API
-void DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
-                       void *Dest,
-                       const void *Src,
-                       size_t Count);
+DPCTLSyclEventRef DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                                    void *Dest,
+                                    const void *Src,
+                                    size_t Count);
 
 /*!
- * @brief C-API wrapper for ``sycl::queue::prefetch``, the function waits on an
- * event till the prefetch operation completes.
+ * @brief C-API wrapper for ``sycl::queue::prefetch``.
  *
  * @param    QRef           An opaque pointer to the ``sycl::queue``.
  * @param    Ptr            An USM pointer to memory.
  * @param    Count          A number of bytes to prefetch.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::prefetch`` function.
  * @ingroup QueueInterface
  */
 DPCTL_API
-void DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
-                         const void *Ptr,
-                         size_t Count);
+DPCTLSyclEventRef DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
+                                      const void *Ptr,
+                                      size_t Count);
 
 /*!
- * @brief C-API wrapper for sycl::queue::mem_advise, the function waits on an
- * event till the operation completes.
+ * @brief C-API wrapper for ``sycl::queue::mem_advise``.
  *
  * @param    QRef           An opaque pointer to the ``sycl::queue``.
  * @param    Ptr            An USM pointer to memory.
  * @param    Count          A number of bytes to prefetch.
  * @param    Advice         Device-defined advice for the specified allocation.
  *                          A value of 0 reverts the advice for Ptr to the
  *                          default behavior.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::mem_advise`` function.
  * @ingroup QueueInterface
  */
 DPCTL_API
-void DPCTLQueue_MemAdvise(__dpctl_keep DPCTLSyclQueueRef QRef,
-                          const void *Ptr,
-                          size_t Count,
-                          int Advice);
+DPCTLSyclEventRef DPCTLQueue_MemAdvise(__dpctl_keep DPCTLSyclQueueRef QRef,
+                                       const void *Ptr,
+                                       size_t Count,
+                                       int Advice);
 
 /*!
  * @brief C-API wrapper for sycl::queue::is_in_order that indicates whether

dpctl-capi/source/dpctl_sycl_queue_interface.cpp

@@ -484,39 +484,74 @@ void DPCTLQueue_Wait(__dpctl_keep DPCTLSyclQueueRef QRef)
     }
 }
 
-void DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
-                       void *Dest,
-                       const void *Src,
-                       size_t Count)
+DPCTLSyclEventRef DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                                    void *Dest,
+                                    const void *Src,
+                                    size_t Count)
 {
     auto Q = unwrap(QRef);
     if (Q) {
-        auto event = Q->memcpy(Dest, Src, Count);
-        event.wait();
+        sycl::event ev;
+        try {
+            ev = Q->memcpy(Dest, Src, Count);
+        } catch (const sycl::runtime_error &re) {
+            // todo: log error
+            std::cerr << re.what() << '\n';
+            return nullptr;
+        }
+        return wrap(new event(ev));
+    }
+    else {
+        // todo: log error
+        std::cerr << "QRef passed to memcpy was NULL" << '\n';
+        return nullptr;
     }
 }
 
-void DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
-                         const void *Ptr,
-                         size_t Count)
+DPCTLSyclEventRef DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
+                                      const void *Ptr,
+                                      size_t Count)
 {
     auto Q = unwrap(QRef);
     if (Q) {
-        auto event = Q->prefetch(Ptr, Count);
-        event.wait();
+        sycl::event ev;
+        try {
+            ev = Q->prefetch(Ptr, Count);
+        } catch (sycl::runtime_error &re) {
+            // todo: log error
+            std::cerr << re.what() << '\n';
+            return nullptr;
+        }
+        return wrap(new event(ev));
+    }
+    else {
+        // todo: log error
+        std::cerr << "QRef passed to prefetch was NULL" << '\n';
+        return nullptr;
     }
 }
 
-void DPCTLQueue_MemAdvise(__dpctl_keep DPCTLSyclQueueRef QRef,
-                          const void *Ptr,
-                          size_t Count,
-                          int Advice)
+DPCTLSyclEventRef DPCTLQueue_MemAdvise(__dpctl_keep DPCTLSyclQueueRef QRef,
+                                       const void *Ptr,
+                                       size_t Count,
+                                       int Advice)
 {
     auto Q = unwrap(QRef);
     if (Q) {
-        auto event =
-            Q->mem_advise(Ptr, Count, static_cast<pi_mem_advice>(Advice));
-        event.wait();
+        sycl::event ev;
+        try {
+            ev = Q->mem_advise(Ptr, Count, static_cast<pi_mem_advice>(Advice));
+        } catch (const sycl::runtime_error &re) {
+            // todo: log error
+            std::cerr << re.what() << '\n';
+            return nullptr;
+        }
+        return wrap(new event(ev));
+    }
+    else {
+        // todo: log error
+        std::cerr << "QRef passed to prefetch was NULL" << '\n';
+        return nullptr;
     }
 }
 

dpctl-capi/tests/test_sycl_queue_interface.cpp

@@ -327,6 +327,24 @@ TEST(TestDPCTLSyclQueueInterface, CheckPropertyHandling)
     EXPECT_NO_FATAL_FAILURE(DPCTLDeviceSelector_Delete(DSRef));
 }
 
+TEST(TestDPCTLSyclQueueInterface, CheckMemOpsZeroQRef)
+{
+    DPCTLSyclQueueRef QRef = nullptr;
+    void *p1 = nullptr;
+    void *p2 = nullptr;
+    size_t n_bytes = 0;
+    DPCTLSyclEventRef ERef = nullptr;
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Memcpy(QRef, p1, p2, n_bytes));
+    ASSERT_FALSE(bool(ERef));
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Prefetch(QRef, p1, n_bytes));
+    ASSERT_FALSE(bool(ERef));
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_MemAdvise(QRef, p1, n_bytes, 0));
+    ASSERT_FALSE(bool(ERef));
+}
+
 TEST_P(TestDPCTLQueueMemberFunctions, CheckGetBackend)
 {
     auto q = unwrap(QRef);
@@ -364,6 +382,31 @@ TEST_P(TestDPCTLQueueMemberFunctions, CheckGetDevice)
     EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(D));
 }
 
+TEST_P(TestDPCTLQueueMemberFunctions, CheckMemOpsNullPtr)
+{
+    void *p1 = nullptr;
+    void *p2 = nullptr;
+    size_t n_bytes = 256;
+    DPCTLSyclEventRef ERef = nullptr;
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Memcpy(QRef, p1, p2, n_bytes));
+    ASSERT_FALSE(bool(ERef));
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_Prefetch(QRef, p1, n_bytes));
+    if (ERef) {
+        ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(ERef));
+        ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(ERef));
+        ERef = nullptr;
+    }
+
+    ASSERT_NO_FATAL_FAILURE(ERef = DPCTLQueue_MemAdvise(QRef, p1, n_bytes, 0));
+    if (ERef) {
+        ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Wait(ERef));
+        ASSERT_NO_FATAL_FAILURE(DPCTLEvent_Delete(ERef));
+        ERef = nullptr;
+    }
+}
+
 INSTANTIATE_TEST_SUITE_P(
     DPCTLQueueMemberFuncTests,
     TestDPCTLQueueMemberFunctions,

dpctl-capi/tests/test_sycl_usm_interface.cpp

@@ -56,12 +56,23 @@ void common_test_body(size_t nbytes,
     auto QueueDev = DPCTLQueue_GetDevice(Q);
     EXPECT_TRUE(DPCTLDevice_AreEq(Dev, QueueDev));
 
-    EXPECT_NO_FATAL_FAILURE(DPCTLQueue_Prefetch(Q, Ptr, nbytes));
-    EXPECT_NO_FATAL_FAILURE(DPCTLQueue_MemAdvise(Q, Ptr, nbytes, 0));
-
+    DPCTLSyclEventRef E1Ref = nullptr, E2Ref = nullptr, E3Ref = nullptr;
+    EXPECT_NO_FATAL_FAILURE(E1Ref = DPCTLQueue_Prefetch(Q, Ptr, nbytes));
+    EXPECT_TRUE(E1Ref != nullptr);
+    EXPECT_NO_FATAL_FAILURE(E2Ref = DPCTLQueue_MemAdvise(Q, Ptr, nbytes, 0));
+    EXPECT_TRUE(E2Ref != nullptr);
+
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Wait(E1Ref));
+    DPCTLEvent_Delete(E1Ref);
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Wait(E2Ref));
+    DPCTLEvent_Delete(E2Ref);
     try {
         unsigned short *host_ptr = new unsigned short[nbytes];
-        EXPECT_NO_FATAL_FAILURE(DPCTLQueue_Memcpy(Q, host_ptr, Ptr, nbytes));
+        EXPECT_NO_FATAL_FAILURE(
+            E3Ref = DPCTLQueue_Memcpy(Q, host_ptr, Ptr, nbytes));
+        EXPECT_TRUE(E3Ref != nullptr);
+        EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Wait(E3Ref));
+        DPCTLEvent_Delete(E3Ref);
         delete[] host_ptr;
     } catch (std::bad_alloc const &ba) {
         // pass

dpctl/_backend.pxd

@@ -355,16 +355,16 @@ cdef extern from "dpctl_sycl_queue_interface.h":
         const DPCTLSyclEventRef *DepEvents,
         size_t NDepEvents)
     cdef void DPCTLQueue_Wait(const DPCTLSyclQueueRef QRef)
-    cdef void DPCTLQueue_Memcpy(
+    cdef DPCTLSyclEventRef DPCTLQueue_Memcpy(
         const DPCTLSyclQueueRef Q,
         void *Dest,
         const void *Src,
         size_t Count)
-    cdef void DPCTLQueue_Prefetch(
+    cdef DPCTLSyclEventRef DPCTLQueue_Prefetch(
         const DPCTLSyclQueueRef Q,
         const void *Src,
         size_t Count)
-    cdef void DPCTLQueue_MemAdvise(
+    cdef DPCTLSyclEventRef DPCTLQueue_MemAdvise(
         const DPCTLSyclQueueRef Q,
         const void *Src,
         size_t Count,

dpctl/_sycl_queue.pyx

@@ -30,6 +30,8 @@ from ._backend cimport (  # noqa: E211
     DPCTLDevice_Delete,
     DPCTLDeviceMgr_GetCachedContext,
     DPCTLDeviceSelector_Delete,
+    DPCTLEvent_Delete,
+    DPCTLEvent_Wait,
     DPCTLFilterSelector_Create,
     DPCTLQueue_AreEq,
     DPCTLQueue_Copy,
@@ -812,6 +814,7 @@ cdef class SyclQueue(_SyclQueue):
     cpdef memcpy(self, dest, src, size_t count):
         cdef void *c_dest
         cdef void *c_src
+        cdef DPCTLSyclEventRef ERef = NULL
 
         if isinstance(dest, _Memory):
             c_dest = <void*>(<_Memory>dest).memory_ptr
@@ -823,10 +826,17 @@ cdef class SyclQueue(_SyclQueue):
         else:
             raise TypeError("Parameter `src` should have type _Memory.")
 
-        DPCTLQueue_Memcpy(self._queue_ref, c_dest, c_src, count)
+        ERef = DPCTLQueue_Memcpy(self._queue_ref, c_dest, c_src, count)
+        if (ERef is NULL):
+            raise RuntimeError(
+                "SyclQueue.memcpy operation encountered an error"
+            )
+        DPCTLEvent_Wait(ERef)
+        DPCTLEvent_Delete(ERef)
 
     cpdef prefetch(self, mem, size_t count=0):
         cdef void *ptr
+        cdef DPCTLSyclEventRef ERef = NULL
 
         if isinstance(mem, _Memory):
             ptr = <void*>(<_Memory>mem).memory_ptr
@@ -836,10 +846,17 @@ cdef class SyclQueue(_SyclQueue):
         if (count <=0 or count > self.nbytes):
             count = self.nbytes
 
-        DPCTLQueue_Prefetch(self._queue_ref, ptr, count)
+        ERef = DPCTLQueue_Prefetch(self._queue_ref, ptr, count)
+        if (ERef is NULL):
+            raise RuntimeError(
+                "SyclQueue.prefetch encountered an error"
+            )
+        DPCTLEvent_Wait(ERef)
+        DPCTLEvent_Delete(ERef)
 
     cpdef mem_advise(self, mem, size_t count, int advice):
         cdef void *ptr
+        cdef DPCTLSyclEventRef ERef = NULL
 
         if isinstance(mem, _Memory):
             ptr = <void*>(<_Memory>mem).memory_ptr
@@ -849,7 +866,13 @@ cdef class SyclQueue(_SyclQueue):
         if (count <=0 or count > self.nbytes):
             count = self.nbytes
 
-        DPCTLQueue_MemAdvise(self._queue_ref, ptr, count, advice)
+        ERef = DPCTLQueue_MemAdvise(self._queue_ref, ptr, count, advice)
+        if (ERef is NULL):
+            raise RuntimeError(
+                "SyclQueue.mem_advise operation encountered an error"
+            )
+        DPCTLEvent_Wait(ERef)
+        DPCTLEvent_Delete(ERef)
 
     @property
     def is_in_order(self):