SyclEvent deleter no longer waits

Added _host_task_utils.hpp file that provides function to
submit host task to decrement reference count for each
Python object passed to one gated by passed event references.

Host task is submitted to sycl::queue behind the given QRef.

SyclQueue.submit changes the mechanism by which it ensures that
Python objects (and memory they hold) that the kernel works on
are not GC collected before the kernel completes its execution.

We used to store reference to arguments in the SyclEvent object
returned by SyclQueue.submit. To avoid GC-ing kernel arguments
before kernel completes execution, deleter of SyclEvent has to
call `SyclEvent.wait`.

With this PR `SyclQueue.submit` increments reference counts of
arguments after submission, then schedules a host task on the same
queue, dependent on the kernel event, which aquires GIL and
decrements reference counts.

After this change the SyclEven.__dealloc__ no longer has to wait.
It also no longer needs to store args attribute. This is now
deperecated and will be removed.

Author: oleksandr-pavlyk

dpctl/CMakeLists.txt

@@ -161,6 +161,8 @@ foreach(_cy_file ${_cython_sources})
     build_dpctl_ext(${_trgt} ${_cy_file} "dpctl")
 endforeach()
 
+target_include_directories(_sycl_queue PRIVATE ${CMAKE_CURRENT_SOURCE_DIR})
+
 add_subdirectory(program)
 add_subdirectory(memory)
 add_subdirectory(tensor)

dpctl/_host_task_util.hpp

@@ -0,0 +1,42 @@
+#include "Python.h"
+#include "syclinterface/dpctl_data_types.h"
+#include <CL/sycl.hpp>
+
+int async_dec_ref(DPCTLSyclQueueRef QRef,
+                  PyObject **obj_array,
+                  size_t obj_array_size,
+                  DPCTLSyclEventRef *ERefs,
+                  size_t nERefs)
+{
+
+    sycl::queue *q = reinterpret_cast<sycl::queue *>(QRef);
+
+    std::vector<PyObject *> obj_vec;
+    obj_vec.reserve(obj_array_size);
+    for (size_t obj_id = 0; obj_id < obj_array_size; ++obj_id) {
+        obj_vec.push_back(obj_array[obj_id]);
+    }
+
+    try {
+        q->submit([&](sycl::handler &cgh) {
+            for (size_t ev_id = 0; ev_id < nERefs; ++ev_id) {
+                cgh.depends_on(
+                    *(reinterpret_cast<sycl::event *>(ERefs[ev_id])));
+            }
+            cgh.host_task([obj_array_size, obj_vec]() {
+                {
+                    PyGILState_STATE gstate;
+                    gstate = PyGILState_Ensure();
+                    for (size_t i = 0; i < obj_array_size; ++i) {
+                        Py_DECREF(obj_vec[i]);
+                    }
+                    PyGILState_Release(gstate);
+                }
+            });
+        });
+    } catch (const std::exception &e) {
+        return 1;
+    }
+
+    return 0;
+}

dpctl/_sycl_event.pyx

@@ -98,7 +98,7 @@ cdef class _SyclEvent:
 
     def __dealloc__(self):
         if (self._event_ref):
-            with nogil: DPCTLEvent_Wait(self._event_ref)
+            # with nogil: DPCTLEvent_Wait(self._event_ref)
             DPCTLEvent_Delete(self._event_ref)
         self._event_ref = NULL
         self.args = None

dpctl/_sycl_queue.pyx

@@ -65,11 +65,17 @@ import ctypes
 from .enum_types import backend_type
 
 from cpython cimport pycapsule
+from cpython.ref cimport Py_INCREF, PyObject
 from libc.stdlib cimport free, malloc
 
 import collections.abc
 import logging
 
+
+cdef extern from "_host_task_util.hpp":
+    int async_dec_ref(DPCTLSyclQueueRef, PyObject **, size_t, DPCTLSyclEventRef *, size_t) nogil
+
+
 __all__ = [
     "SyclQueue",
     "SyclKernelInvalidRangeError",
@@ -714,12 +720,14 @@ cdef class SyclQueue(_SyclQueue):
         cdef _arg_data_type *kargty = NULL
         cdef DPCTLSyclEventRef *depEvents = NULL
         cdef DPCTLSyclEventRef Eref = NULL
-        cdef int ret
+        cdef int ret = 0
         cdef size_t gRange[3]
         cdef size_t lRange[3]
         cdef size_t nGS = len(gS)
         cdef size_t nLS = len(lS) if lS is not None else 0
         cdef size_t nDE = len(dEvents) if dEvents is not None else 0
+        cdef PyObject **arg_objects = NULL
+        cdef ssize_t i = 0
 
         # Allocate the arrays to be sent to DPCTLQueue_Submit
         kargs = <void**>malloc(len(args) * sizeof(void*))
@@ -820,8 +828,18 @@ cdef class SyclQueue(_SyclQueue):
             raise SyclKernelSubmitError(
                 "Kernel submission to Sycl queue failed."
             )
-
-        return SyclEvent._create(Eref, args)
+        # increment reference counts to each argument
+        arg_objects = <PyObject **>malloc(len(args) * sizeof(PyObject *))
+        for i in range(len(args)):
+            arg_objects[i] = <PyObject *>(args[i])
+            Py_INCREF(<object> arg_objects[i])
+
+        # schedule decrement
+        async_dec_ref(self.get_queue_ref(), arg_objects, len(args), &Eref, 1)
+        # free memory
+        free(arg_objects)
+
+        return SyclEvent._create(Eref, [])
 
     cpdef void wait(self):
         with nogil: DPCTLQueue_Wait(self._queue_ref)