From b6710edfb5a1656a04d7168d6776c3d3554b1648 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 13 Aug 2021 11:03:11 -0500 Subject: [PATCH 1/2] Implemented DPCTLQueue_SubmitBarrier, DPCTLQueue_SubmitBarrierForEvents --- .../include/dpctl_sycl_queue_interface.h | 28 ++++++++++ .../source/dpctl_sycl_queue_interface.cpp | 41 +++++++++++++++ dpctl-capi/tests/test_sycl_queue_submit.cpp | 52 +++++++++++++++++++ 3 files changed, 121 insertions(+) diff --git a/dpctl-capi/include/dpctl_sycl_queue_interface.h b/dpctl-capi/include/dpctl_sycl_queue_interface.h index 08cfae264d..f4e4d0b7e6 100644 --- a/dpctl-capi/include/dpctl_sycl_queue_interface.h +++ b/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 diff --git a/dpctl-capi/source/dpctl_sycl_queue_interface.cpp b/dpctl-capi/source/dpctl_sycl_queue_interface.cpp index 3ccaca7511..6a3ffb15fd 100644 --- a/dpctl-capi/source/dpctl_sycl_queue_interface.cpp +++ b/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); +} diff --git a/dpctl-capi/tests/test_sycl_queue_submit.cpp b/dpctl-capi/tests/test_sycl_queue_submit.cpp index ca9fb3f233..04441347f9 100644 --- a/dpctl-capi/tests/test_sycl_queue_submit.cpp +++ b/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 From 484330b92ea2b542e347fdcc2ad449d72782e785 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 13 Aug 2021 12:23:31 -0500 Subject: [PATCH 2/2] Implemented SyclQueue.submit_barrier --- dpctl/_backend.pxd | 6 ++++++ dpctl/_sycl_queue.pxd | 1 + dpctl/_sycl_queue.pyx | 39 ++++++++++++++++++++++++++++++++++ dpctl/tests/test_sycl_queue.py | 13 ++++++++++++ 4 files changed, 59 insertions(+) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 14917afe78..d514d064b9 100644 --- a/dpctl/_backend.pxd +++ b/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": diff --git a/dpctl/_sycl_queue.pxd b/dpctl/_sycl_queue.pxd index 3b5eeb9c1b..4fd7058781 100644 --- a/dpctl/_sycl_queue.pxd +++ b/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=*) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 41586fd3b5..198584280d 100644 --- a/dpctl/_sycl_queue.pyx +++ b/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( 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 = ( + malloc(nDE*sizeof(DPCTLSyclEventRef)) + ) + if not depEvents: + raise MemoryError() + else: + for idx, de in enumerate(dependent_events): + depEvents[idx] = (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, []) diff --git a/dpctl/tests/test_sycl_queue.py b/dpctl/tests/test_sycl_queue.py index 9d3b1008a7..0422368629 100644 --- a/dpctl/tests/test_sycl_queue.py +++ b/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()