Initial async kernel support

Author: ZzEeKkAa

numba_dpex/core/runtime/CMakeLists.txt

@@ -109,7 +109,7 @@ python_add_library(${PROJECT_NAME} MODULE ${SOURCES})
 
 # Add SYCL to target, this must come after python_add_library()
 # FIXME: sources incompatible with sycl include?
-# add_sycl_to_target(TARGET ${PROJECT_NAME} SOURCES ${KERNEL_SOURCES})
+add_sycl_to_target(TARGET ${PROJECT_NAME} SOURCES ${KERNEL_SOURCES})
 
 # Link the DPCTLSyclInterface library to target
 target_link_libraries(${PROJECT_NAME} PRIVATE DPCTLSyclInterface)

numba_dpex/core/runtime/_dbg_printer.h

@@ -13,7 +13,7 @@
 
 /* Debugging facilities - enabled at compile-time */
 /* #undef NDEBUG */
-#if 0
+#if 1
 #include <stdio.h>
 #define DPEXRT_DEBUG(X)                                                        \
     {                                                                          \

numba_dpex/core/runtime/_dpexrt_python.c

@@ -24,6 +24,7 @@
 #include "_queuestruct.h"
 #include "_usmarraystruct.h"
 
+#include "experimental/experimental.h"
 #include "numba/core/runtime/nrt_external.h"
 
 // forward declarations
@@ -1490,6 +1491,8 @@ static PyObject *build_c_helpers_dict(void)
                  &DPEXRT_sycl_event_from_python);
     _declpointer("DPEXRT_sycl_event_to_python", &DPEXRT_sycl_event_to_python);
     _declpointer("DPEXRT_sycl_event_init", &DPEXRT_sycl_event_init);
+    _declpointer("DPEXRT_nrt_acuire_meminfo_and_schedule_release",
+                 &DPEXRT_nrt_acuire_meminfo_and_schedule_release);
 
 #undef _declpointer
     return dct;
@@ -1557,6 +1560,9 @@ MOD_INIT(_dpexrt_python)
                        PyLong_FromVoidPtr(&DPEXRT_MemInfo_alloc));
     PyModule_AddObject(m, "DPEXRT_MemInfo_fill",
                        PyLong_FromVoidPtr(&DPEXRT_MemInfo_fill));
+    PyModule_AddObject(
+        m, "DPEXRT_nrt_acuire_meminfo_and_schedule_release",
+        PyLong_FromVoidPtr(&DPEXRT_nrt_acuire_meminfo_and_schedule_release));
     PyModule_AddObject(m, "c_helpers", build_c_helpers_dict());
     return MOD_SUCCESS_VAL(m);
 }

numba_dpex/core/runtime/context.py

@@ -433,3 +433,41 @@ def submit_ndrange(
         )
 
         return ret
+
+    def acuire_meminfo_and_schedule_release(
+        self, builder: llvmir.IRBuilder, args
+    ):
+        """Inserts LLVM IR to call nrt_acuire_meminfo_and_schedule_release.
+
+        DPCTLSyclEventRef
+        DPEXRT_nrt_acuire_meminfo_and_schedule_release(
+            NRT_api_functions *nrt,
+            DPCTLSyclQueueRef QRef,
+            NRT_MemInfo **meminfo_array,
+            size_t meminfo_array_size,
+            DPCTLSyclEventRef *depERefs,
+            size_t nDepERefs,
+            int *status,
+        );
+
+        """
+        mod = builder.module
+
+        func_ty = llvmir.FunctionType(
+            cgutils.voidptr_t,
+            [
+                cgutils.voidptr_t,
+                cgutils.voidptr_t,
+                cgutils.voidptr_t.as_pointer(),
+                llvmir.IntType(64),
+                cgutils.voidptr_t,
+                llvmir.IntType(64),
+                llvmir.IntType(64).as_pointer(),
+            ],
+        )
+        fn = cgutils.get_or_insert_function(
+            mod, func_ty, "DPEXRT_nrt_acuire_meminfo_and_schedule_release"
+        )
+        ret = builder.call(fn, args)
+
+        return ret

numba_dpex/core/runtime/experimental/experimental.h

@@ -0,0 +1,24 @@
+#ifndef _EXPERIMENTAL_H_
+#define _EXPERIMENTAL_H_
+
+#include "dpctl_capi.h"
+#include "numba/core/runtime/nrt_external.h"
+
+#ifdef __cplusplus
+extern "C"
+{
+#endif
+    DPCTLSyclEventRef DPEXRT_nrt_acuire_meminfo_and_schedule_release(
+        NRT_api_functions *nrt,
+        DPCTLSyclQueueRef QRef,
+        NRT_MemInfo **meminfo_array,
+        size_t meminfo_array_size,
+        // DPCTLSyclEventRef *depERefs,
+        DPCTLSyclEventRef depERef,
+        size_t nDepERefs,
+        int *status);
+#ifdef __cplusplus
+}
+#endif
+
+#endif /* _EXPERIMENTAL_H_ */

numba_dpex/core/runtime/experimental/sycl_release_nrt.cpp

@@ -0,0 +1,69 @@
+#include "experimental.h"
+
+#include "_dbg_printer.h"
+#include "syclinterface/dpctl_sycl_type_casters.hpp"
+#include <CL/sycl.hpp>
+
+extern "C"
+{
+    DPCTLSyclEventRef DPEXRT_nrt_acuire_meminfo_and_schedule_release(
+        NRT_api_functions *nrt,
+        DPCTLSyclQueueRef QRef,
+        NRT_MemInfo **meminfo_array,
+        size_t meminfo_array_size,
+        // DPCTLSyclEventRef *depERefs,
+        DPCTLSyclEventRef depERef,
+        size_t nDepERefs,
+        int *status)
+    {
+        DPEXRT_DEBUG(drt_debug_print(
+                         "DPEXRT-DEBUG: scheduling nrt meminfo release.\n"););
+
+        using dpctl::syclinterface::unwrap;
+        using dpctl::syclinterface::wrap;
+
+        sycl::queue *q = unwrap<sycl::queue>(QRef);
+
+        std::vector<NRT_MemInfo *> meminfo_vec(
+            meminfo_array, meminfo_array + meminfo_array_size);
+
+        for (size_t i = 0; i < meminfo_array_size; ++i) {
+            nrt->acquire(meminfo_vec[i]);
+        }
+
+        DPEXRT_DEBUG(drt_debug_print("DPEXRT-DEBUG: acquired meminfo.\n"););
+
+        try {
+            sycl::event ht_ev = q->submit([&](sycl::handler &cgh) {
+                for (size_t ev_id = 0; ev_id < nDepERefs; ++ev_id) {
+                    // cgh.depends_on(*(unwrap<sycl::event>(depERefs[ev_id])));
+                    cgh.depends_on(*(unwrap<sycl::event>(depERef)));
+                }
+                cgh.host_task([meminfo_array_size, meminfo_vec, nrt]() {
+                    for (size_t i = 0; i < meminfo_array_size; ++i) {
+                        nrt->release(meminfo_vec[i]);
+                        DPEXRT_DEBUG(
+                            drt_debug_print("DPEXRT-DEBUG: released meminfo "
+                                            "from host_task.\n"););
+                    }
+                });
+            });
+
+            constexpr int result_ok = 0;
+
+            *status = result_ok;
+            auto e_ptr = new sycl::event(ht_ev);
+            return wrap<sycl::event>(e_ptr);
+        } catch (const std::exception &e) {
+            constexpr int result_std_exception = 1;
+
+            *status = result_std_exception;
+            return nullptr;
+        }
+
+        constexpr int result_other_abnormal = 2;
+
+        *status = result_other_abnormal;
+        return nullptr;
+    }
+}