Feature/submit barrier

dpctl-capi/include/dpctl_sycl_queue_interface.h

@@ -343,4 +343,32 @@ bool DPCTLQueue_IsInOrder(__dpctl_keep const DPCTLSyclQueueRef QRef);
 DPCTL_API
 size_t DPCTLQueue_Hash(__dpctl_keep const DPCTLSyclQueueRef QRef);
 
+/*!
+ * @brief C-API wraper for ``sycl::queue::submit_barrier()``.
+ *
+ * @param    QRef    An opaque pointer to the ``sycl::queue``.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::submit_barrier()`` function.
+ */
+DPCTL_API
+DPCTLSyclEventRef
+DPCTLQueue_SubmitBarrier(__dpctl_keep const DPCTLSyclQueueRef QRef);
+
+/*!
+ * @brief C-API wraper for ``sycl::queue::submit_barrier(event_vector)``.
+ *
+ * @param    QRef    An opaque pointer to the ``sycl::queue``.
+ * @param    DepEvents     List of dependent DPCTLSyclEventRef objects (events)
+ *                         for the barrier. We call ``sycl::handler.depends_on``
+ *                         for each of the provided events.
+ * @param    NDepEvents    Size of the DepEvents list.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::submit_barrier()`` function.
+ */
+DPCTL_API
+DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents(
+    __dpctl_keep const DPCTLSyclQueueRef QRef,
+    __dpctl_keep const DPCTLSyclEventRef *DepEvents,
+    size_t NDepEvents);
+
 DPCTL_C_EXTERN_C_END

dpctl-capi/source/dpctl_sycl_queue_interface.cpp

@@ -533,3 +533,44 @@ size_t DPCTLQueue_Hash(__dpctl_keep const DPCTLSyclQueueRef QRef)
         return 0;
     }
 }
+
+__dpctl_give DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents(
+    __dpctl_keep const DPCTLSyclQueueRef QRef,
+    __dpctl_keep const DPCTLSyclEventRef *DepEvents,
+    size_t NDepEvents)
+{
+    auto Q = unwrap(QRef);
+    event e;
+    if (Q) {
+        try {
+            e = Q->submit([&](handler &cgh) {
+                // Depend on any event that was specified by the caller.
+                if (NDepEvents)
+                    for (auto i = 0ul; i < NDepEvents; ++i)
+                        cgh.depends_on(*unwrap(DepEvents[i]));
+
+                cgh.barrier();
+            });
+        } catch (runtime_error &re) {
+            // \todo fix error handling
+            std::cerr << re.what() << '\n';
+            return nullptr;
+        } catch (std::runtime_error &sre) {
+            std::cerr << sre.what() << '\n';
+            return nullptr;
+        }
+
+        return wrap(new event(e));
+    }
+    else {
+        // todo: log error
+        std::cerr << "Argument QRef is null" << '\n';
+        return nullptr;
+    }
+}
+
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_SubmitBarrier(__dpctl_keep const DPCTLSyclQueueRef QRef)
+{
+    return DPCTLQueue_SubmitBarrierForEvents(QRef, nullptr, 0);
+}

dpctl-capi/tests/test_sycl_queue_submit.cpp

@@ -356,4 +356,56 @@ TEST_F(TestQueueSubmitNDRange, ChkSubmitNDRangeDouble)
     EXPECT_TRUE(worked);
 }
 
+struct TestQueueSubmitBarrier : public ::testing::Test
+{
+    DPCTLSyclQueueRef QRef = nullptr;
+
+    TestQueueSubmitBarrier()
+    {
+        DPCTLSyclDeviceSelectorRef DSRef = nullptr;
+        DPCTLSyclDeviceRef DRef = nullptr;
+
+        EXPECT_NO_FATAL_FAILURE(DSRef = DPCTLDefaultSelector_Create());
+        EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef));
+        EXPECT_NO_FATAL_FAILURE(QRef = DPCTLQueue_CreateForDevice(
+                                    DRef, nullptr, DPCTL_DEFAULT_PROPERTY));
+        EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef));
+        EXPECT_NO_FATAL_FAILURE(DPCTLDeviceSelector_Delete(DSRef));
+    }
+    ~TestQueueSubmitBarrier()
+    {
+        EXPECT_NO_FATAL_FAILURE(DPCTLQueue_Delete(QRef));
+    }
+};
+
+TEST_F(TestQueueSubmitBarrier, ChkSubmitBarrier)
+{
+    DPCTLSyclEventRef ERef = nullptr;
+
+    ASSERT_TRUE(QRef != nullptr);
+    EXPECT_NO_FATAL_FAILURE(ERef = DPCTLQueue_SubmitBarrier(QRef));
+    ASSERT_TRUE(ERef != nullptr);
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Wait(ERef));
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Delete(ERef));
+}
+
+TEST_F(TestQueueSubmitBarrier, ChkSubmitBarrierWithEvents)
+{
+    DPCTLSyclEventRef ERef = nullptr;
+    DPCTLSyclEventRef DepsERefs[2] = {nullptr, nullptr};
+
+    EXPECT_NO_FATAL_FAILURE(DepsERefs[0] = DPCTLEvent_Create());
+    EXPECT_NO_FATAL_FAILURE(DepsERefs[1] = DPCTLEvent_Create());
+
+    ASSERT_TRUE(QRef != nullptr);
+    EXPECT_NO_FATAL_FAILURE(
+        ERef = DPCTLQueue_SubmitBarrierForEvents(QRef, DepsERefs, 2));
+
+    ASSERT_TRUE(ERef != nullptr);
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Wait(ERef));
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Delete(ERef));
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Delete(DepsERefs[0]));
+    EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Delete(DepsERefs[1]));
+}
+
 #endif

dpctl/_backend.pxd

@@ -347,6 +347,12 @@ cdef extern from "dpctl_sycl_queue_interface.h":
         size_t Count,
         int Advice)
     cdef bool DPCTLQueue_IsInOrder(const DPCTLSyclQueueRef QRef)
+    cdef DPCTLSyclEventRef DPCTLQueue_SubmitBarrier(
+        const DPCTLSyclQueueRef QRef)
+    cdef DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents(
+        const DPCTLSyclQueueRef QRef,
+        const DPCTLSyclEventRef *DepEvents,
+        size_t NDepEvents)
 
 
 cdef extern from "dpctl_sycl_queue_manager.h":

dpctl/_sycl_queue.pxd

@@ -81,3 +81,4 @@ cdef public class SyclQueue (_SyclQueue) [object PySyclQueueObject, type PySyclQ
     cpdef memcpy(self, dest, src, size_t count)
     cpdef prefetch(self, ptr, size_t count=*)
     cpdef mem_advise(self, ptr, size_t count, int mem)
+    cpdef SyclEvent submit_barrier(self, dependent_events=*)

dpctl/_sycl_queue.pyx

@@ -43,6 +43,7 @@ from ._backend cimport (  # noqa: E211
     DPCTLQueue_MemAdvise,
     DPCTLQueue_Memcpy,
     DPCTLQueue_Prefetch,
+    DPCTLQueue_SubmitBarrierForEvents,
     DPCTLQueue_SubmitNDRange,
     DPCTLQueue_SubmitRange,
     DPCTLQueue_Wait,
@@ -63,6 +64,7 @@ from .enum_types import backend_type
 from cpython cimport pycapsule
 from libc.stdlib cimport free, malloc
 
+import collections.abc
 import logging
 
 __all__ = [
@@ -879,3 +881,40 @@ cdef class SyclQueue(_SyclQueue):
         return pycapsule.PyCapsule_New(
             <void *>QRef, "SyclQueueRef", &_queue_capsule_deleter
         )
+
+    cpdef SyclEvent submit_barrier(self, dependent_events=None):
+        """
+        Submits a barrier to the queue.
+        """
+        cdef DPCTLSyclEventRef *depEvents = NULL
+        cdef DPCTLSyclEventRef ERef = NULL
+        cdef size_t nDE = 0
+        # Create the array of dependent events if any
+        if (dependent_events is None or
+            (isinstance(dependent_events, collections.abc.Sequence) and
+             all([type(de) is SyclEvent for de in dependent_events]))):
+            nDE = 0 if dependent_events is None else len(dependent_events)
+        else:
+            raise TypeError(
+                "dependent_events must either None, or a sequence of "
+                ":class:`dpctl.SyclEvent` objects")
+        if nDE > 0:
+            depEvents = (
+                <DPCTLSyclEventRef*>malloc(nDE*sizeof(DPCTLSyclEventRef))
+            )
+            if not depEvents:
+                raise MemoryError()
+            else:
+                for idx, de in enumerate(dependent_events):
+                    depEvents[idx] = (<SyclEvent>de).get_event_ref()
+
+        ERef = DPCTLQueue_SubmitBarrierForEvents(
+            self.get_queue_ref(), depEvents, nDE)
+        if (depEvents is not NULL):
+            free(depEvents)
+        if ERef is NULL:
+            raise SyclKernelSubmitError(
+                "Barrier submission to Sycl queue failed."
+            )
+
+        return SyclEvent._create(ERef, [])

dpctl/tests/test_sycl_queue.py

@@ -381,3 +381,16 @@ def test_hashing_of_queue():
     """
     queue_dict = {dpctl.SyclQueue(): "default_queue"}
     assert queue_dict
+
+
+def test_queue_submit_barrier(valid_filter):
+    try:
+        q = dpctl.SyclQueue(valid_filter)
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Failed to create device with supported filter")
+    ev1 = q.submit_barrier()
+    ev2 = q.submit_barrier()
+    ev3 = q.submit_barrier([ev1, ev2])
+    ev3.wait()
+    ev1.wait()
+    ev2.wait()