diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 7d2b5af54b..03ec7f21d1 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -229,6 +229,15 @@ cdef extern from "dpctl_sycl_event_interface.h": cdef void DPCTLEvent_Delete(DPCTLSyclEventRef ERef) cdef _event_status_type DPCTLEvent_GetCommandExecutionStatus(DPCTLSyclEventRef ERef) cdef _backend_type DPCTLEvent_GetBackend(DPCTLSyclEventRef ERef) + cdef struct DPCTLEventVector + ctypedef DPCTLEventVector *DPCTLEventVectorRef + cdef void DPCTLEventVector_Delete(DPCTLEventVectorRef EVRef) + cdef size_t DPCTLEventVector_Size(DPCTLEventVectorRef EVRef) + cdef DPCTLSyclEventRef DPCTLEventVector_GetAt( + DPCTLEventVectorRef EVRef, + size_t index) + cdef DPCTLEventVectorRef DPCTLEvent_GetWaitList( + DPCTLSyclEventRef ERef) cdef extern from "dpctl_sycl_kernel_interface.h": diff --git a/dpctl/_sycl_event.pyx b/dpctl/_sycl_event.pyx index 6a767fd69d..580d4c2ae0 100644 --- a/dpctl/_sycl_event.pyx +++ b/dpctl/_sycl_event.pyx @@ -31,7 +31,12 @@ from ._backend cimport ( # noqa: E211 DPCTLEvent_Delete, DPCTLEvent_GetBackend, DPCTLEvent_GetCommandExecutionStatus, + DPCTLEvent_GetWaitList, DPCTLEvent_Wait, + DPCTLEventVector_Delete, + DPCTLEventVector_GetAt, + DPCTLEventVector_Size, + DPCTLEventVectorRef, DPCTLSyclEventRef, _backend_type, _event_status_type, @@ -240,3 +245,20 @@ cdef class SyclEventRaw(_SyclEventRaw): return backend_type.cuda else: raise ValueError("Unknown backend type.") + + def get_wait_list(self): + cdef DPCTLEventVectorRef EVRef = DPCTLEvent_GetWaitList( + self.get_event_ref() + ) + cdef size_t num_events + cdef size_t i + cdef DPCTLSyclEventRef ERef + if (EVRef is NULL): + raise ValueError("Internal error: NULL event vector encountered") + num_events = DPCTLEventVector_Size(EVRef) + events = [] + for i in range(num_events): + ERef = DPCTLEventVector_GetAt(EVRef, i) + events.append(SyclEventRaw._create(ERef)) + DPCTLEventVector_Delete(EVRef) + return events diff --git a/dpctl/tests/test_sycl_event.py b/dpctl/tests/test_sycl_event.py index 7bdbf8fe11..f298405743 100644 --- a/dpctl/tests/test_sycl_event.py +++ b/dpctl/tests/test_sycl_event.py @@ -87,3 +87,50 @@ def test_backend(): dpctl.SyclEventRaw().backend except ValueError: pytest.fail("Failed to get backend from event") + + +@pytest.mark.skip(reason="event::get_wait_list() method returns wrong result") +def test_get_wait_list(): + if has_cpu(): + oclSrc = " \ + kernel void add_k(global float* a) { \ + size_t index = get_global_id(0); \ + a[index] = a[index] + 1; \ + } \ + kernel void sqrt_k(global float* a) { \ + size_t index = get_global_id(0); \ + a[index] = sqrt(a[index]); \ + } \ + kernel void sin_k(global float* a) { \ + size_t index = get_global_id(0); \ + a[index] = sin(a[index]); \ + }" + q = dpctl.SyclQueue("opencl:cpu") + prog = dpctl_prog.create_program_from_source(q, oclSrc) + addKernel = prog.get_sycl_kernel("add_k") + sqrtKernel = prog.get_sycl_kernel("sqrt_k") + sinKernel = prog.get_sycl_kernel("sin_k") + + bufBytes = 1024 * np.dtype("f").itemsize + abuf = dpctl_mem.MemoryUSMShared(bufBytes, queue=q) + a = np.ndarray((1024), buffer=abuf, dtype="f") + a[:] = np.arange(1024) + args = [] + + args.append(a.base) + r = [1024] + ev_1 = q.submit(addKernel, args, r) + ev_2 = q.submit(sqrtKernel, args, r, dEvents=[ev_1]) + ev_3 = q.submit(sinKernel, args, r, dEvents=[ev_2]) + + ev_raw = dpctl.SyclEventRaw(ev_3) + + try: + wait_list = ev_raw.get_wait_list() + except ValueError: + pytest.fail( + "Failed to get a list of waiting events from SyclEventRaw" + ) + assert len(wait_list) + else: + pytest.skip("No OpenCL CPU queues available")