C API for SyclProgram and SyclKernel classes

Author: ndgrigorian

dpctl/_backend.pxd

@@ -263,6 +263,7 @@ cdef extern from "syclinterface/dpctl_sycl_event_interface.h":
 cdef extern from "syclinterface/dpctl_sycl_kernel_interface.h":
     cdef size_t DPCTLKernel_GetNumArgs(const DPCTLSyclKernelRef KRef)
     cdef void DPCTLKernel_Delete(DPCTLSyclKernelRef KRef)
+    cdef DPCTLSyclKernelRef DPCTLKernel_Copy(const DPCTLSyclKernelRef KRef)
     cdef size_t DPCTLKernel_GetWorkGroupSize(const DPCTLSyclKernelRef KRef)
     cdef size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple(const DPCTLSyclKernelRef KRef)
     cdef size_t DPCTLKernel_GetPrivateMemSize(const DPCTLSyclKernelRef KRef)
@@ -341,6 +342,7 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
     cdef bool DPCTLKernelBundle_HasKernel(DPCTLSyclKernelBundleRef KBRef,
                                      const char *KernelName)
     cdef void DPCTLKernelBundle_Delete(DPCTLSyclKernelBundleRef KBRef)
+    cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(const DPCTLSyclKernelBundleRef KBRef)
 
 
 cdef extern from "syclinterface/dpctl_sycl_queue_interface.h":

dpctl/apis/include/dpctl4pybind11.hpp

@@ -71,6 +71,8 @@ struct dpctl_capi
     PyTypeObject *PyMemoryUSMSharedType_;
     PyTypeObject *PyMemoryUSMHostType_;
     PyTypeObject *PyUSMArrayType_;
+    PyTypeObject *PySyclProgramType_;
+    PyTypeObject *PySyclKernelType_;
 
     DPCTLSyclDeviceRef (*SyclDevice_GetDeviceRef_)(PySyclDeviceObject *);
     PySyclDeviceObject *(*SyclDevice_Make_)(DPCTLSyclDeviceRef);
@@ -94,6 +96,13 @@ struct dpctl_capi
                               DPCTLSyclQueueRef,
                               PyObject *);
 
+    // program
+    DPCTLSyclKernelRef (*SyclKernel_GetKernelRef_)(PySyclKernelObject *);
+    PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef);
+
+    DPCTLSyclKernelBundleRef (*SyclProgram_GetKernelBundleRef_)(PySyclProgramObject *);
+    PySyclProgramObject *(*SyclProgram_Make_)(DPCTLSyclKernelBundleRef);
+
     // tensor
     char *(*UsmNDArray_GetData_)(PyUSMArrayObject *);
     int (*UsmNDArray_GetNDim_)(PyUSMArrayObject *);
@@ -131,6 +140,14 @@ struct dpctl_capi
     {
         return PyObject_TypeCheck(obj, PySyclQueueType_) != 0;
     }
+    bool PySyclKernel_Check_(PyObject *obj) const
+    {
+        return PyObject_TypeCheck(obj, PySyclKernelType_) != 0;
+    }
+    bool PySyclProgram_Check_(PyObject *obj) const
+    {
+        return PyObject_TypeCheck(obj, PySyclProgramType_) != 0;
+    }
 
     ~dpctl_capi(){};
 
@@ -174,6 +191,8 @@ struct dpctl_capi
     std::shared_ptr<py::object> default_usm_memory;
     std::shared_ptr<py::object> default_usm_ndarray;
     std::shared_ptr<py::object> as_usm_memory;
+    std::shared_ptr<py::object> default_sycl_kernel;
+    std::shared_ptr<py::object> default_sycl_program;
 
     dpctl_capi()
         : default_sycl_queue{}, default_usm_memory{}, default_usm_ndarray{},
@@ -201,6 +220,8 @@ struct dpctl_capi
         this->PyMemoryUSMSharedType_ = &PyMemoryUSMSharedType;
         this->PyMemoryUSMHostType_ = &PyMemoryUSMHostType;
         this->PyUSMArrayType_ = &PyUSMArrayType;
+        this->PySyclProgramType_ = &PySyclProgramType;
+        this->PySyclKernelType_ = &PySyclKernelType;
 
         // SyclDevice API
         this->SyclDevice_GetDeviceRef_ = SyclDevice_GetDeviceRef;
@@ -225,6 +246,10 @@ struct dpctl_capi
         this->Memory_GetNumBytes_ = Memory_GetNumBytes;
         this->Memory_Make_ = Memory_Make;
 
+        // dpctl.program API
+        this->SyclKernel_Make_ = SyclKernel_Make;
+        this->SyclProgram_Make_ = SyclProgram_Make;
+
         // dpctl.tensor.usm_ndarray API
         this->UsmNDArray_GetData_ = UsmNDArray_GetData;
         this->UsmNDArray_GetNDim_ = UsmNDArray_GetNDim;
@@ -506,6 +531,76 @@ template <> struct type_caster<sycl::event>
 
     DPCTL_TYPE_CASTER(sycl::event, _("dpctl.SyclEvent"));
 };
+
+/* This type caster associates ``sycl::kernel`` C++ class with
+ * :class:`dpctl.program.SyclKernel` for the purposes of generation of
+ * Python bindings by pybind11.
+ */
+template <> struct type_caster<sycl::kernel>
+{
+public:
+    bool load(handle src, bool)
+    {
+        PyObject *source = src.ptr();
+        auto &api = ::dpctl::detail::dpctl_capi::get();
+        if (api.PySyclKernel_Check_(source)) {
+            DPCTLSyclKernelRef KRef = api.SyclKernel_GetKernelRef_(
+                reinterpret_cast<PySyclKernelObject *>(source));
+            value = std::make_unique<sycl::kernel>(
+                *(reinterpret_cast<sycl::kernel *>(KRef)));
+            return true;
+        }
+        else {
+            throw py::type_error(
+                "Input is of unexpected type, expected dpctl.program.SyclKernel");
+        }
+    }
+
+    static handle cast(sycl::kernel src, return_value_policy, handle)
+    {
+        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto tmp =
+            api.SyclKernel_Make_(reinterpret_cast<DPCTLSyclKernelRef>(&src));
+        return handle(reinterpret_cast<PyObject *>(tmp));
+    }
+
+    DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.program.SyclKernel"));
+};
+
+/* This type caster associates ``sycl::kernel_bundle<sycl::bundle_state::executable>`` C++ class with
+ * :class:`dpctl.program.SyclProgram` for the purposes of generation of
+ * Python bindings by pybind11.
+ */
+template <> struct type_caster<sycl::kernel_bundle<sycl::bundle_state::executable>>
+{
+public:
+    bool load(handle src, bool)
+    {
+        PyObject *source = src.ptr();
+        auto &api = ::dpctl::detail::dpctl_capi::get();
+        if (api.PySyclProgram_Check_(source)) {
+            DPCTLSyclKernelBundleRef KBRef = api.SyclProgram_GetKernelBundleRef_(
+                reinterpret_cast<PySyclProgramObject *>(source));
+            value = std::make_unique<sycl::kernel_bundle<sycl::bundle_state::executable>>(
+                *(reinterpret_cast<sycl::kernel_bundle<sycl::bundle_state::executable> *>(KBRef)));
+            return true;
+        }
+        else {
+            throw py::type_error(
+                "Input is of unexpected type, expected dpctl.SyclEvent");
+        }
+    }
+
+    static handle cast(sycl::kernel_bundle<sycl::bundle_state::executable> src, return_value_policy, handle)
+    {
+        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto tmp =
+            api.SyclProgram_Make_(reinterpret_cast<DPCTLSyclKernelBundleRef>(&src));
+        return handle(reinterpret_cast<PyObject *>(tmp));
+    }
+
+    DPCTL_TYPE_CASTER(sycl::kernel_bundle<sycl::bundle_state::executable>, _("dpctl.program.SyclProgram"));
+};
 } // namespace detail
 } // namespace pybind11
 

dpctl/apis/include/dpctl_capi.h

@@ -40,6 +40,10 @@
 #include "../memory/_memory_api.h"
 #include "../tensor/_usmarray.h"
 #include "../tensor/_usmarray_api.h"
+#include "../program/_program.h"
+#include "../program/_program_api.h"
+
+
 // clang-format on
 
 /*
@@ -59,5 +63,6 @@ static inline void import_dpctl(void)
     import_dpctl___sycl_queue();
     import_dpctl__memory___memory();
     import_dpctl__tensor___usmarray();
+    import_dpctl__program___program();
     return;
 }

dpctl/program/_program.pxd

@@ -28,7 +28,7 @@ from .._sycl_device cimport SyclDevice
 from .._sycl_queue cimport SyclQueue
 
 
-cdef class SyclKernel:
+cdef api class SyclKernel [object PySyclKernelObject, type PySyclKernelType]:
     ''' Wraps a sycl::kernel object created from an OpenCL interoperability
         kernel.
     '''
@@ -40,7 +40,7 @@ cdef class SyclKernel:
     cdef SyclKernel _create (DPCTLSyclKernelRef kref, str name)
 
 
-cdef class SyclProgram:
+cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]:
     ''' Wraps a sycl::kernel_bundle<sycl::bundle_state::executable> object created by
         using SYCL interoperability layer for OpenCL and Level-Zero backends.
 
@@ -59,3 +59,4 @@ cdef class SyclProgram:
 cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*)
 cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL,
                                  unicode copts=*)
+

dpctl/program/_program.pyx

@@ -31,6 +31,7 @@ from libc.stdint cimport uint32_t
 from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLCString_Delete,
     DPCTLKernel_Delete,
+    DPCTLKernel_Copy,
     DPCTLKernel_GetCompileNumSubGroups,
     DPCTLKernel_GetCompileSubGroupSize,
     DPCTLKernel_GetMaxNumSubGroups,
@@ -41,6 +42,7 @@ from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLKernelBundle_CreateFromOCLSource,
     DPCTLKernelBundle_CreateFromSpirv,
     DPCTLKernelBundle_Delete,
+    DPCTLKernelBundle_Copy,
     DPCTLKernelBundle_GetKernel,
     DPCTLKernelBundle_HasKernel,
     DPCTLSyclContextRef,
@@ -164,6 +166,19 @@ cdef class SyclKernel:
         cdef size_t n = DPCTLKernel_GetCompileSubGroupSize(self._kernel_ref)
         return n
 
+cdef api DPCTLSyclKernelRef SyclKernel_GetKernelRef(SyclKernel ker):
+    """ C-API function to access opaque kernel reference from
+    Python object of type :class:`dpctl.program.SyclKernel`.
+    """
+    return ker.get_kernel_ref()
+
+cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef):
+    """
+    C-API function to create :class:`dpctl.program.SyclKernel`
+    instance from opaque sycl kernel reference.
+    """
+    cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef)
+    return SyclKernel._create(copied_KRef, "foo")
 
 cdef class SyclProgram:
     """ Wraps a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object
@@ -290,3 +305,17 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
         raise SyclProgramCompilationError()
 
     return SyclProgram._create(KBref)
+
+cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef(SyclProgram pro):
+    """ C-API function to access opaque kernel bundle reference from
+    Python object of type :class:`dpctl.program.SyclKernel`.
+    """
+    return pro.get_program_ref()
+
+cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef):
+    """
+    C-API function to create :class:`dpctl.program.SyclProgram`
+    instance from opaque sycl kernel bundle reference.
+    """
+    cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef)
+    return SyclProgram._create(copied_KBRef)

libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h

@@ -117,4 +117,16 @@ bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef,
 DPCTL_API
 void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef);
 
+/*!
+ * @brief Returns a copy of the DPCTLSyclKernelBundleRef object.
+ *
+ * @param    KBRef           DPCTLSyclKernelBundleRef object to be copied.
+ * @return   A new DPCTLSyclKernelBundleRef created by copying the passed in
+ * DPCTLSyclKernelBundleRef object.
+ * @ingroup KernelBundleInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclKernelBundleRef
+DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef);
+
 DPCTL_C_EXTERN_C_END

libsyclinterface/include/dpctl_sycl_kernel_interface.h

@@ -62,6 +62,18 @@ size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef);
 DPCTL_API
 void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef);
 
+/*!
+ * @brief Returns a copy of the DPCTLSyclKernelRef object.
+ *
+ * @param    KRef           DPCTLSyclKernelRef object to be copied.
+ * @return   A new DPCTLSyclKernelRef created by copying the passed in
+ * DPCTLSyclKernelRef object.
+ * @ingroup KernelInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclKernelRef
+DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef);
+
 /*!
  * !brief Wrapper around
  * `kernel::get_info<info::kernel_device_specific::work_group_size>()`.

libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp

@@ -740,3 +740,21 @@ void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef)
 {
     delete unwrap(KBRef);
 }
+
+__dpctl_give DPCTLSyclKernelBundleRef
+DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef)
+{
+    auto Bundle = unwrap(KBRef);
+    if (!Bundle) {
+        error_handler("Cannot copy DPCTLSyclKernelBundleRef as input is a nullptr",
+                      __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+    try {
+        auto CopiedBundle = new kernel_bundle<bundle_state::executable>(*Bundle);
+        return wrap(CopiedBundle);
+    } catch (std::exception const &e) {
+        error_handler(e, __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+}

libsyclinterface/source/dpctl_sycl_kernel_interface.cpp

@@ -59,6 +59,24 @@ void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef)
     delete unwrap(KRef);
 }
 
+__dpctl_give DPCTLSyclKernelRef
+DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef)
+{
+    auto Kernel = unwrap(KRef);
+    if (!Kernel) {
+        error_handler("Cannot copy DPCTLSyclKernelRef as input is a nullptr",
+                      __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+    try {
+        auto CopiedKernel = new kernel(*Kernel);
+        return wrap(CopiedKernel);
+    } catch (std::exception const &e) {
+        error_handler(e, __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+}
+
 size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
 {
     if (!KRef) {