diff --git a/MANIFEST.in b/MANIFEST.in index 7a1dcbf027..2cb0991412 100644 --- a/MANIFEST.in +++ b/MANIFEST.in @@ -12,6 +12,8 @@ include dpctl/_sycl_event.h include dpctl/_sycl_event_api.h include dpctl/memory/_memory.h include dpctl/memory/_memory_api.h +include dpctl/program/_program.h +include dpctl/program/_program_api.h include dpctl/tensor/_usmarray.h include dpctl/tensor/_usmarray_api.h recursive-include dpctl/tensor/include * diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index fdd19a6652..f4ce41f6a2 100644 --- a/dpctl/_backend.pxd +++ b/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": diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 5b40f222ce..f2a5dc11e9 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -54,9 +54,9 @@ constexpr int platform_typeid_lookup(int I, Ints... Is) : platform_typeid_lookup(Is...); } -struct dpctl_capi +class dpctl_capi { - +public: // dpctl type objects PyTypeObject *Py_SyclDeviceType_; PyTypeObject *PySyclDeviceType_; @@ -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,14 @@ struct dpctl_capi DPCTLSyclQueueRef, PyObject *); + // program + DPCTLSyclKernelRef (*SyclKernel_GetKernelRef_)(PySyclKernelObject *); + PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef, const char *); + + DPCTLSyclKernelBundleRef (*SyclProgram_GetKernelBundleRef_)( + PySyclProgramObject *); + PySyclProgramObject *(*SyclProgram_Make_)(DPCTLSyclKernelBundleRef); + // tensor char *(*UsmNDArray_GetData_)(PyUSMArrayObject *); int (*UsmNDArray_GetNDim_)(PyUSMArrayObject *); @@ -131,6 +141,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(){}; @@ -142,19 +160,19 @@ struct dpctl_capi py::object default_sycl_queue_pyobj() { - return *default_sycl_queue; + return *default_sycl_queue_; } py::object default_usm_memory_pyobj() { - return *default_usm_memory; + return *default_usm_memory_; } py::object default_usm_ndarray_pyobj() { - return *default_usm_ndarray; + return *default_usm_ndarray_; } py::object as_usm_memory_pyobj() { - return *as_usm_memory; + return *as_usm_memory_; } private: @@ -170,14 +188,14 @@ struct dpctl_capi } }; - std::shared_ptr default_sycl_queue; - std::shared_ptr default_usm_memory; - std::shared_ptr default_usm_ndarray; - std::shared_ptr as_usm_memory; + std::shared_ptr default_sycl_queue_; + std::shared_ptr default_usm_memory_; + std::shared_ptr default_usm_ndarray_; + std::shared_ptr as_usm_memory_; dpctl_capi() - : default_sycl_queue{}, default_usm_memory{}, default_usm_ndarray{}, - as_usm_memory{} + : default_sycl_queue_{}, default_usm_memory_{}, default_usm_ndarray_{}, + as_usm_memory_{} { // Import Cython-generated C-API for dpctl // This imports python modules and initializes @@ -201,6 +219,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 +245,12 @@ struct dpctl_capi this->Memory_GetNumBytes_ = Memory_GetNumBytes; this->Memory_Make_ = Memory_Make; + // dpctl.program API + this->SyclKernel_GetKernelRef_ = SyclKernel_GetKernelRef; + this->SyclKernel_Make_ = SyclKernel_Make; + this->SyclProgram_GetKernelBundleRef_ = SyclProgram_GetKernelBundleRef; + this->SyclProgram_Make_ = SyclProgram_Make; + // dpctl.tensor.usm_ndarray API this->UsmNDArray_GetData_ = UsmNDArray_GetData; this->UsmNDArray_GetNDim_ = UsmNDArray_GetNDim; @@ -284,18 +310,18 @@ struct dpctl_capi py::object py_sycl_queue = py::reinterpret_steal( reinterpret_cast(py_q_tmp)); - default_sycl_queue = std::shared_ptr( + default_sycl_queue_ = std::shared_ptr( new py::object(py_sycl_queue), Deleter{}); py::module_ mod_memory = py::module_::import("dpctl.memory"); py::object py_as_usm_memory = mod_memory.attr("as_usm_memory"); - as_usm_memory = std::shared_ptr( + as_usm_memory_ = std::shared_ptr( new py::object{py_as_usm_memory}, Deleter{}); auto mem_kl = mod_memory.attr("MemoryUSMHost"); py::object py_default_usm_memory = mem_kl(1, py::arg("queue") = py_sycl_queue); - default_usm_memory = std::shared_ptr( + default_usm_memory_ = std::shared_ptr( new py::object{py_default_usm_memory}, Deleter{}); py::module_ mod_usmarray = @@ -306,7 +332,7 @@ struct dpctl_capi tensor_kl(py::tuple(), py::arg("dtype") = py::str("u1"), py::arg("buffer") = py_default_usm_memory); - default_usm_ndarray = std::shared_ptr( + default_usm_ndarray_ = std::shared_ptr( new py::object{py_default_usm_ndarray}, Deleter{}); } @@ -377,7 +403,7 @@ template <> struct type_caster bool load(handle src, bool) { PyObject *source = src.ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); if (api.PySyclQueue_Check_(source)) { DPCTLSyclQueueRef QRef = api.SyclQueue_GetQueueRef_( reinterpret_cast(source)); @@ -393,7 +419,7 @@ template <> struct type_caster static handle cast(sycl::queue src, return_value_policy, handle) { - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); auto tmp = api.SyclQueue_Make_(reinterpret_cast(&src)); return handle(reinterpret_cast(tmp)); @@ -412,7 +438,7 @@ template <> struct type_caster bool load(handle src, bool) { PyObject *source = src.ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); if (api.PySyclDevice_Check_(source)) { DPCTLSyclDeviceRef DRef = api.SyclDevice_GetDeviceRef_( reinterpret_cast(source)); @@ -428,7 +454,7 @@ template <> struct type_caster static handle cast(sycl::device src, return_value_policy, handle) { - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); auto tmp = api.SyclDevice_Make_(reinterpret_cast(&src)); return handle(reinterpret_cast(tmp)); @@ -447,7 +473,7 @@ template <> struct type_caster bool load(handle src, bool) { PyObject *source = src.ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); if (api.PySyclContext_Check_(source)) { DPCTLSyclContextRef CRef = api.SyclContext_GetContextRef_( reinterpret_cast(source)); @@ -463,7 +489,7 @@ template <> struct type_caster static handle cast(sycl::context src, return_value_policy, handle) { - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); auto tmp = api.SyclContext_Make_(reinterpret_cast(&src)); return handle(reinterpret_cast(tmp)); @@ -482,7 +508,7 @@ template <> struct type_caster bool load(handle src, bool) { PyObject *source = src.ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); if (api.PySyclEvent_Check_(source)) { DPCTLSyclEventRef ERef = api.SyclEvent_GetEventRef_( reinterpret_cast(source)); @@ -498,7 +524,7 @@ template <> struct type_caster static handle cast(sycl::event src, return_value_policy, handle) { - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); auto tmp = api.SyclEvent_Make_(reinterpret_cast(&src)); return handle(reinterpret_cast(tmp)); @@ -506,6 +532,86 @@ template <> struct type_caster 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 +{ +public: + bool load(handle src, bool) + { + PyObject *source = src.ptr(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); + if (api.PySyclKernel_Check_(source)) { + DPCTLSyclKernelRef KRef = api.SyclKernel_GetKernelRef_( + reinterpret_cast(source)); + value = std::make_unique( + *(reinterpret_cast(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 const &api = ::dpctl::detail::dpctl_capi::get(); + auto tmp = + api.SyclKernel_Make_(reinterpret_cast(&src), + "dpctl4pybind11_kernel"); + return handle(reinterpret_cast(tmp)); + } + + DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.program.SyclKernel")); +}; + +/* This type caster associates + * ``sycl::kernel_bundle`` C++ class with + * :class:`dpctl.program.SyclProgram` for the purposes of generation of + * Python bindings by pybind11. + */ +template <> +struct type_caster> +{ +public: + bool load(handle src, bool) + { + PyObject *source = src.ptr(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); + if (api.PySyclProgram_Check_(source)) { + DPCTLSyclKernelBundleRef KBRef = + api.SyclProgram_GetKernelBundleRef_( + reinterpret_cast(source)); + value = std::make_unique< + sycl::kernel_bundle>( + *(reinterpret_cast< + sycl::kernel_bundle *>( + KBRef))); + return true; + } + else { + throw py::type_error("Input is of unexpected type, expected " + "dpctl.program.SyclProgram"); + } + } + + static handle cast(sycl::kernel_bundle src, + return_value_policy, + handle) + { + auto const &api = ::dpctl::detail::dpctl_capi::get(); + auto tmp = api.SyclProgram_Make_( + reinterpret_cast(&src)); + return handle(reinterpret_cast(tmp)); + } + + DPCTL_TYPE_CASTER(sycl::kernel_bundle, + _("dpctl.program.SyclProgram")); +}; } // namespace detail } // namespace pybind11 @@ -544,7 +650,7 @@ class usm_memory : public py::object sycl::queue get_queue() const { Py_MemoryObject *mem_obj = reinterpret_cast(m_ptr); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); DPCTLSyclQueueRef QRef = api.Memory_GetQueueRef_(mem_obj); sycl::queue *obj_q = reinterpret_cast(QRef); return *obj_q; @@ -553,14 +659,14 @@ class usm_memory : public py::object char *get_pointer() const { Py_MemoryObject *mem_obj = reinterpret_cast(m_ptr); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); DPCTLSyclUSMRef MRef = api.Memory_GetUsmPointer_(mem_obj); return reinterpret_cast(MRef); } size_t get_nbytes() const { - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); Py_MemoryObject *mem_obj = reinterpret_cast(m_ptr); return api.Memory_GetNumBytes_(mem_obj); } @@ -663,7 +769,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return api.UsmNDArray_GetData_(raw_ar); } @@ -676,7 +782,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return api.UsmNDArray_GetNDim_(raw_ar); } @@ -684,7 +790,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return api.UsmNDArray_GetShape_(raw_ar); } @@ -698,7 +804,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return api.UsmNDArray_GetStrides_(raw_ar); } @@ -706,7 +812,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); int ndim = api.UsmNDArray_GetNDim_(raw_ar); const py::ssize_t *shape = api.UsmNDArray_GetShape_(raw_ar); @@ -723,7 +829,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); int nd = api.UsmNDArray_GetNDim_(raw_ar); const py::ssize_t *shape = api.UsmNDArray_GetShape_(raw_ar); const py::ssize_t *strides = api.UsmNDArray_GetStrides_(raw_ar); @@ -757,7 +863,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); DPCTLSyclQueueRef QRef = api.UsmNDArray_GetQueueRef_(raw_ar); return *(reinterpret_cast(QRef)); } @@ -766,7 +872,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return api.UsmNDArray_GetTypenum_(raw_ar); } @@ -774,7 +880,7 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return api.UsmNDArray_GetFlags_(raw_ar); } @@ -782,28 +888,28 @@ class usm_ndarray : public py::object { PyUSMArrayObject *raw_ar = this->usm_array_ptr(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return api.UsmNDArray_GetElementSize_(raw_ar); } bool is_c_contiguous() const { int flags = this->get_flags(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return static_cast(flags & api.USM_ARRAY_C_CONTIGUOUS_); } bool is_f_contiguous() const { int flags = this->get_flags(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return static_cast(flags & api.USM_ARRAY_F_CONTIGUOUS_); } bool is_writable() const { int flags = this->get_flags(); - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); return static_cast(flags & api.USM_ARRAY_WRITABLE_); } diff --git a/dpctl/apis/include/dpctl_capi.h b/dpctl/apis/include/dpctl_capi.h index d6c104581a..9715e42b38 100644 --- a/dpctl/apis/include/dpctl_capi.h +++ b/dpctl/apis/include/dpctl_capi.h @@ -40,6 +40,9 @@ #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 +62,6 @@ static inline void import_dpctl(void) import_dpctl___sycl_queue(); import_dpctl__memory___memory(); import_dpctl__tensor___usmarray(); + import_dpctl__program___program(); return; } diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 2133a7a667..86c338aff7 100644 --- a/dpctl/program/_program.pxd +++ b/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 object created by using SYCL interoperability layer for OpenCL and Level-Zero backends. diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 26a9c67a7e..82ff39de56 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -30,6 +30,7 @@ from libc.stdint cimport uint32_t from dpctl._backend cimport ( # noqa: E211, E402; DPCTLCString_Delete, + DPCTLKernel_Copy, DPCTLKernel_Delete, DPCTLKernel_GetCompileNumSubGroups, DPCTLKernel_GetCompileSubGroupSize, @@ -38,6 +39,7 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_GetPreferredWorkGroupSizeMultiple, DPCTLKernel_GetPrivateMemSize, DPCTLKernel_GetWorkGroupSize, + DPCTLKernelBundle_Copy, DPCTLKernelBundle_CreateFromOCLSource, DPCTLKernelBundle_CreateFromSpirv, DPCTLKernelBundle_Delete, @@ -165,6 +167,25 @@ cdef class SyclKernel: 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, const char *name): + """ + C-API function to create :class:`dpctl.program.SyclKernel` + instance from opaque sycl kernel reference. + """ + cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef) + if (name is NULL): + return SyclKernel._create(copied_KRef, "default_name") + else: + return SyclKernel._create(copied_KRef, name.decode("utf-8")) + + cdef class SyclProgram: """ Wraps a ``sycl::kernel_bundle`` object created using SYCL interoperability layer with underlying backends. Only the @@ -290,3 +311,19 @@ 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) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index ac55dce972..1cbd9114e4 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -31,6 +31,123 @@ def get_spirv_abspath(fn): return spirv_file +def _check_cpython_api_SyclProgram_GetKernelBundleRef(sycl_prog): + """Checks Cython-generated C-API function + `SyclProgram_GetKernelBundleRef` defined in _program.pyx""" + import ctypes + import sys + + assert type(sycl_prog) is dpctl_prog.SyclProgram + mod = sys.modules[sycl_prog.__class__.__module__] + # get capsule storing SyclProgram_GetKernelBundleRef function ptr + kb_ref_fn_cap = mod.__pyx_capi__["SyclProgram_GetKernelBundleRef"] + # construct Python callable to invoke "SyclProgram_GetKernelBundleRef" + cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer + cap_ptr_fn.restype = ctypes.c_void_p + cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p] + kb_ref_fn_ptr = cap_ptr_fn( + kb_ref_fn_cap, + b"DPCTLSyclKernelBundleRef (struct PySyclProgramObject *)", + ) + # PYFUNCTYPE(result_type, *arg_types) + callable_maker = ctypes.PYFUNCTYPE(ctypes.c_void_p, ctypes.py_object) + get_kernel_bundle_ref_fn = callable_maker(kb_ref_fn_ptr) + + r2 = sycl_prog.addressof_ref() + r1 = get_kernel_bundle_ref_fn(sycl_prog) + assert r1 == r2 + + +def _check_cpython_api_SyclProgram_Make(sycl_prog): + """Checks Cython-generated C-API function + `SyclProgram_Make` defined in _program.pyx""" + import ctypes + import sys + + assert type(sycl_prog) is dpctl_prog.SyclProgram + mod = sys.modules[sycl_prog.__class__.__module__] + # get capsule storing SyclProgram_Make function ptr + make_prog_fn_cap = mod.__pyx_capi__["SyclProgram_Make"] + # construct Python callable to invoke "SyclProgram_Make" + cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer + cap_ptr_fn.restype = ctypes.c_void_p + cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p] + make_prog_fn_ptr = cap_ptr_fn( + make_prog_fn_cap, + b"struct PySyclProgramObject *(DPCTLSyclKernelBundleRef)", + ) + # PYFUNCTYPE(result_type, *arg_types) + callable_maker = ctypes.PYFUNCTYPE(ctypes.py_object, ctypes.c_void_p) + make_prog_fn = callable_maker(make_prog_fn_ptr) + + p2 = make_prog_fn(sycl_prog.addressof_ref()) + assert p2.has_sycl_kernel("add") + assert p2.has_sycl_kernel("axpy") + + +def _check_cpython_api_SyclKernel_GetKernelRef(krn): + """Checks Cython-generated C-API function + `SyclKernel_GetKernelRef` defined in _program.pyx""" + import ctypes + import sys + + assert type(krn) is dpctl_prog.SyclKernel + mod = sys.modules[krn.__class__.__module__] + # get capsule storing SyclKernel_GetKernelRef function ptr + k_ref_fn_cap = mod.__pyx_capi__["SyclKernel_GetKernelRef"] + # construct Python callable to invoke "SyclKernel_GetKernelRef" + cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer + cap_ptr_fn.restype = ctypes.c_void_p + cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p] + k_ref_fn_ptr = cap_ptr_fn( + k_ref_fn_cap, b"DPCTLSyclKernelRef (struct PySyclKernelObject *)" + ) + # PYFUNCTYPE(result_type, *arg_types) + callable_maker = ctypes.PYFUNCTYPE(ctypes.c_void_p, ctypes.py_object) + get_kernel_ref_fn = callable_maker(k_ref_fn_ptr) + + r2 = krn.addressof_ref() + r1 = get_kernel_ref_fn(krn) + assert r1 == r2 + + +def _check_cpython_api_SyclKernel_Make(krn): + """Checks Cython-generated C-API function + `SyclKernel_Make` defined in _program.pyx""" + import ctypes + import sys + + assert type(krn) is dpctl_prog.SyclKernel + mod = sys.modules[krn.__class__.__module__] + # get capsule storing SyclKernel_Make function ptr + k_make_fn_cap = mod.__pyx_capi__["SyclKernel_Make"] + # construct Python callable to invoke "SyclKernel_Make" + cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer + cap_ptr_fn.restype = ctypes.c_void_p + cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p] + k_make_fn_ptr = cap_ptr_fn( + k_make_fn_cap, + b"struct PySyclKernelObject *(DPCTLSyclKernelRef, char const *)", + ) + # PYFUNCTYPE(result_type, *arg_types) + callable_maker = ctypes.PYFUNCTYPE( + ctypes.py_object, ctypes.c_void_p, ctypes.c_void_p + ) + make_kernel_fn = callable_maker(k_make_fn_ptr) + + k2 = make_kernel_fn( + krn.addressof_ref(), bytes(krn.get_function_name(), "utf-8") + ) + assert krn.get_function_name() == k2.get_function_name() + assert krn.get_num_args() == k2.get_num_args() + assert krn.work_group_size == k2.work_group_size + + k3 = make_kernel_fn(krn.addressof_ref(), ctypes.c_void_p(None)) + assert k3.get_function_name() == "default_name" + assert krn.get_num_args() == k3.get_num_args() + assert krn.work_group_size == k3.work_group_size + + def _check_multi_kernel_program(prog): assert type(prog) is dpctl_prog.SyclProgram @@ -49,6 +166,9 @@ def _check_multi_kernel_program(prog): assert type(axpyKernel.addressof_ref()) is int for krn in [addKernel, axpyKernel]: + _check_cpython_api_SyclKernel_GetKernelRef(krn) + _check_cpython_api_SyclKernel_Make(krn) + na = krn.num_args assert na == krn.get_num_args() wgsz = krn.work_group_size @@ -68,6 +188,9 @@ def _check_multi_kernel_program(prog): cmsgsz = krn.compile_sub_group_size assert type(cmsgsz) is int + _check_cpython_api_SyclProgram_GetKernelBundleRef(prog) + _check_cpython_api_SyclProgram_Make(prog) + def test_create_program_from_source_ocl(): oclSrc = " \ diff --git a/dpctl/tests/test_usm_ndarray_ctor.py b/dpctl/tests/test_usm_ndarray_ctor.py index 0531eda402..e00747bfe9 100644 --- a/dpctl/tests/test_usm_ndarray_ctor.py +++ b/dpctl/tests/test_usm_ndarray_ctor.py @@ -825,8 +825,11 @@ def test_astype(): X[:] = np.full((5, 5), 7, dtype="i4") Y = dpt.astype(X, "c8", order="C") assert np.allclose(dpt.to_numpy(Y), np.full((5, 5), 7, dtype="c8")) - Y = dpt.astype(X[::2, ::-1], "f2", order="K") - assert np.allclose(dpt.to_numpy(Y), np.full(Y.shape, 7, dtype="f2")) + if Y.sycl_device.has_aspect_fp16: + Y = dpt.astype(X[::2, ::-1], "f2", order="K") + assert np.allclose(dpt.to_numpy(Y), np.full(Y.shape, 7, dtype="f2")) + Y = dpt.astype(X[::2, ::-1], "f4", order="K") + assert np.allclose(dpt.to_numpy(Y), np.full(Y.shape, 7, dtype="f4")) Y = dpt.astype(X[::2, ::-1], "i4", order="K", copy=False) assert Y.usm_data is X.usm_data diff --git a/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt new file mode 100644 index 0000000000..441a0997c6 --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt @@ -0,0 +1,36 @@ +cmake_minimum_required(VERSION 3.21) + +project(use_queue_device LANGUAGES CXX) + +set(DPCTL_CMAKE_MODULES_PATH "${CMAKE_SOURCE_DIR}/../../../cmake") +set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${DPCTL_CMAKE_MODULES_PATH}) +find_package(IntelDPCPP REQUIRED PATHS ${DPCTL_CMAKE_MODULES_PATH} NO_DEFAULT_PATH) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED True) +set(CMAKE_BUILD_TYPE Debug) + +# Fetch pybind11 +include(FetchContent) +FetchContent_Declare( + pybind11 + URL https://github.com/pybind/pybind11/archive/refs/tags/v2.10.0.tar.gz + URL_HASH SHA256=eacf582fa8f696227988d08cfc46121770823839fe9e301a20fbce67e7cd70ec +) +FetchContent_MakeAvailable(pybind11) + +find_package(PythonExtensions REQUIRED) +find_package(Dpctl REQUIRED) +find_package(NumPy REQUIRED) + +set(py_module_name _use_kernel) +pybind11_add_module(${py_module_name} + MODULE + use_kernel/_example.cpp +) +target_include_directories(${py_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS}) +install(TARGETS ${py_module_name} + DESTINATION use_kernel +) + +set(ignoreMe "${SKBUILD}") diff --git a/examples/pybind11/use_dpctl_sycl_kernel/README.md b/examples/pybind11/use_dpctl_sycl_kernel/README.md new file mode 100644 index 0000000000..77aa57bf6e --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/README.md @@ -0,0 +1,25 @@ +# Usage of dpctl Entities in Pybind11 + +## Description + +This extension demonstrates how you can use dpctl Python types, +such as ``dpctl.SyclQueue`` and ``dpctl.program.SyclKernel``, in +Pybind11 extensions. + + +## Building + +To build the extension, run: +``` +source /opt/intel/oneapi/compiler/latest/env/vars.sh +CXX=icpx python setup.py build_ext --inplace +python -m pytest tests +python example.py +``` + +# Sample output + +``` +(dpctl) [17:25:27 ubuntu_vm use_dpctl_syclkernel]$ python example.py +[ 0 2 4 6 8 10 12 14 16 18 20 22 24] +``` diff --git a/examples/pybind11/use_dpctl_sycl_kernel/example.py b/examples/pybind11/use_dpctl_sycl_kernel/example.py new file mode 100644 index 0000000000..292cdf414d --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/example.py @@ -0,0 +1,47 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2022 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# coding: utf-8 + +import use_kernel as eg + +import dpctl +import dpctl.program as dppr +import dpctl.tensor as dpt + +# create execution queue, targeting default selected device +q = dpctl.SyclQueue() + +# read SPIR-V: a program in Khronos standardized intermediate form +with open("resource/double_it.spv", "br") as fh: + il = fh.read() + +# Build the program for the selected device +pr = dppr.create_program_from_spirv(q, il, "") +assert pr.has_sycl_kernel("double_it") + +# Retrieve the kernel from the problem +krn = pr.get_sycl_kernel("double_it") +assert krn.num_args == 2 + +# Construct the argument, and allocate memory for the result +x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q) +y = dpt.empty_like(x) + +eg.submit_custom_kernel(q, krn, src=x, dst=y) + +# output the result +print(dpt.asnumpy(y)) diff --git a/examples/pybind11/use_dpctl_sycl_kernel/resource/README.md b/examples/pybind11/use_dpctl_sycl_kernel/resource/README.md new file mode 100644 index 0000000000..e83f8296a8 --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/resource/README.md @@ -0,0 +1,8 @@ +# Rebuilding SPIR-V file from source + +```bash +export TOOLS_DIR=$(dirname $(dirname $(which icx)))/bin-llvm +$TOOLS_DIR/clang -cc1 -triple spir double_it.cl -finclude-default-header -flto -emit -llvm-bc -o double_it.bc +$TOOLS_DIR/llvm-spirv double_it.bc -o double_it.spv +rm double_it.bc +``` diff --git a/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl new file mode 100644 index 0000000000..660e4b9b60 --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl @@ -0,0 +1,30 @@ +/* + Data Parallel Control (dpctl) + + Copyright 2020-2022 Intel Corporation + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. + +===----------------------------------------------------------------------=== + + \file + This file implements a sample OpenCL kernel to use in this example. + +===----------------------------------------------------------------------=== +*/ + +__kernel void double_it(__global int *x, __global int *y) { + uint idx = get_global_id(0); + + y[idx] = 2 * x[idx]; +} diff --git a/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv new file mode 100644 index 0000000000..c193a1f02b Binary files /dev/null and b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv differ diff --git a/examples/pybind11/use_dpctl_sycl_kernel/setup.py b/examples/pybind11/use_dpctl_sycl_kernel/setup.py new file mode 100644 index 0000000000..8888a9ff0d --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/setup.py @@ -0,0 +1,26 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2022 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +from skbuild import setup + +setup( + name="use_kernel", + version="0.0.1", + description="an example of SYCL-powered Python package (with pybind11)", + author="Intel Scripting", + license="Apache 2.0", + packages=["use_kernel"], +) diff --git a/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py b/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py new file mode 100644 index 0000000000..a8c96e09c5 --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py @@ -0,0 +1,69 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2022 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# coding: utf-8 + +import os.path + +import numpy as np +import pytest +import use_kernel as uk + +import dpctl +import dpctl.program as dpm +import dpctl.tensor as dpt + + +def _get_spv_path(): + uk_dir = os.path.dirname(os.path.abspath(uk.__file__)) + proj_dir = os.path.dirname(uk_dir) + return os.path.join(proj_dir, "resource", "double_it.spv") + + +def test_spv_file_exists(): + assert os.path.exists(_get_spv_path()) + + +def test_kernel_can_be_found(): + fn = _get_spv_path() + with open(fn, "br") as f: + il = f.read() + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Could not create default queue") + pr = dpm.create_program_from_spirv(q, il, "") + assert pr.has_sycl_kernel("double_it") + + +def test_kernel_submit_through_extension(): + fn = _get_spv_path() + with open(fn, "br") as f: + il = f.read() + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Could not create default queue") + pr = dpm.create_program_from_spirv(q, il, "") + krn = pr.get_sycl_kernel("double_it") + assert krn.num_args == 2 + + x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q) + y = dpt.zeros_like(x) + + uk.submit_custom_kernel(q, krn, x, y, []) + + assert np.array_equal(dpt.asnumpy(y), np.arange(0, 26, step=2, dtype="i4")) diff --git a/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py new file mode 100644 index 0000000000..09ac9ffaf0 --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py @@ -0,0 +1,34 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2022 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# coding: utf-8 + +from ._use_kernel import submit_custom_kernel + +__all__ = [ + "submit_custom_kernel", +] + +__doc__ = """ +Example pybind11 extension demonstrating binding of dpctl entities to +SYCL entities. + +dpctl provides type casters that bind ``sycl::kernel`` to +`dpctl.program.SyclKernel`, ``sycl::device`` to `dpctl.SyclDevice`, etc. + +Use of these type casters simplifies writing of Python extensions and compile +then using SYCL C++ compilers, such as Intel(R) oneAPI DPC++ compiler. +""" diff --git a/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp new file mode 100644 index 0000000000..638db20320 --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp @@ -0,0 +1,92 @@ +//==- _example.cpp - Example of Pybind11 extension working with =---------===// +// dpctl Python objects. +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2022 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This file implements Pybind11-generated extension exposing functions that +/// take dpctl Python objects, such as dpctl.SyclQueue, dpctl.SyclDevice as +/// arguments. +/// +//===----------------------------------------------------------------------===// + +#include "dpctl4pybind11.hpp" +#include +#include +#include +#include +#include + +namespace py = pybind11; + +void submit_custom_kernel(sycl::queue q, + sycl::kernel krn, + dpctl::tensor::usm_ndarray x, + dpctl::tensor::usm_ndarray y, + const std::vector &depends = {}) +{ + if (x.get_ndim() != 1 || !x.is_c_contiguous() || y.get_ndim() != 1 || + !y.is_c_contiguous()) + { + throw py::value_error( + "src and dst arguments must be 1D and contiguous."); + } + + auto const &api = dpctl::detail::dpctl_capi::get(); + if (x.get_typenum() != api.UAR_INT32_ || y.get_typenum() != api.UAR_INT32_) + { + throw py::value_error( + "src and dst arguments must have int32 element data types."); + } + + size_t n_x = x.get_size(); + size_t n_y = y.get_size(); + + if (n_x != n_y) { + throw py::value_error("src and dst arguments must have equal size."); + } + + if (!dpctl::utils::queues_are_compatible(q, {x.get_queue(), y.get_queue()})) + { + throw std::runtime_error( + "Execution queue is not compatible with allocation queues"); + } + + void *x_data = x.get_data(); + void *y_data = y.get_data(); + + sycl::event e = q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + cgh.set_arg(0, x_data); + cgh.set_arg(1, y_data); + cgh.parallel_for(sycl::range<1>(n_x), krn); + }); + + e.wait(); + + return; +} + +PYBIND11_MODULE(_use_kernel, m) +{ + m.def("submit_custom_kernel", &submit_custom_kernel, + "Submit given kernel with arguments (int *, int *) to queue", + py::arg("queue"), py::arg("kernel"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); +} diff --git a/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h index dae11bb3c8..97074a5bc6 100644 --- a/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h +++ b/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 diff --git a/libsyclinterface/include/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_interface.h index 43c5dd7a8c..ad46ecb3f1 100644 --- a/libsyclinterface/include/dpctl_sycl_kernel_interface.h +++ b/libsyclinterface/include/dpctl_sycl_kernel_interface.h @@ -26,6 +26,7 @@ #pragma once +#include "Config/dpctl_config.h" #include "Support/DllExport.h" #include "Support/ExternC.h" #include "Support/MemOwnershipAttrs.h" @@ -62,6 +63,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()`. @@ -117,7 +130,7 @@ DPCTL_API uint32_t DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef); -#if 0 +#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER /*! * !brief Wrapper around * `kernel::get_info()`. diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index a6e0492421..b196264350 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -739,3 +739,23 @@ 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); + return wrap>(CopiedBundle); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } +} diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index cdf009e5b4..95a800a34e 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -25,6 +25,7 @@ //===----------------------------------------------------------------------===// #include "dpctl_sycl_kernel_interface.h" +#include "Config/dpctl_config.h" #include "dpctl_error_handlers.h" #include "dpctl_string_utils.hpp" #include "dpctl_sycl_type_casters.hpp" @@ -57,6 +58,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) { @@ -142,7 +161,7 @@ DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef) return static_cast(v); } -#if 0 +#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER // commented out due to bug in DPC++ runtime, get_info for max_sub_group_size // exported by libsycl has different, not SPEC-compliant signature uint32_t @@ -161,8 +180,9 @@ DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef) __FILE__, __func__, __LINE__); return 0; } - auto v = sycl_kern - ->get_info(devs[0]); + auto v = + sycl_kern->get_info( + devs[0]); return v; } #endif diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index 3560100be6..6383b730a0 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -102,6 +102,27 @@ TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirv) ASSERT_FALSE(DPCTLKernelBundle_HasKernel(KBRef, nullptr)); } +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCopy) +{ + DPCTLSyclKernelBundleRef Copied_KBRef = nullptr; + ASSERT_TRUE(KBRef != nullptr); + + EXPECT_NO_FATAL_FAILURE(Copied_KBRef = DPCTLKernelBundle_Copy(KBRef)); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(Copied_KBRef, "add")); + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(Copied_KBRef, "axpy")); + + EXPECT_NO_FATAL_FAILURE(DPCTLKernelBundle_Delete(Copied_KBRef)); +} + +TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCopyNullArgument) +{ + DPCTLSyclKernelBundleRef Null_KBRef = nullptr; + DPCTLSyclKernelBundleRef Copied_KBRef = nullptr; + + EXPECT_NO_FATAL_FAILURE(Copied_KBRef = DPCTLKernelBundle_Copy(Null_KBRef)); + ASSERT_TRUE(Copied_KBRef == nullptr); +} + TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvNull) { DPCTLSyclContextRef Null_CRef = nullptr; diff --git a/libsyclinterface/tests/test_sycl_kernel_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_interface.cpp index a45eae034f..b679fe09d5 100644 --- a/libsyclinterface/tests/test_sycl_kernel_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_interface.cpp @@ -24,6 +24,7 @@ /// //===----------------------------------------------------------------------===// +#include "Config/dpctl_config.h" #include "dpctl_sycl_context_interface.h" #include "dpctl_sycl_device_interface.h" #include "dpctl_sycl_device_selector_interface.h" @@ -99,11 +100,18 @@ struct TestDPCTLSyclKernelInterface TEST_P(TestDPCTLSyclKernelInterface, CheckGetNumArgs) { - ASSERT_EQ(DPCTLKernel_GetNumArgs(AddKRef), 3ul); ASSERT_EQ(DPCTLKernel_GetNumArgs(AxpyKRef), 4ul); } +TEST_P(TestDPCTLSyclKernelInterface, CheckCopy) +{ + DPCTLSyclKernelRef Copied_KRef = nullptr; + EXPECT_NO_FATAL_FAILURE(Copied_KRef = DPCTLKernel_Copy(AddKRef)); + ASSERT_EQ(DPCTLKernel_GetNumArgs(Copied_KRef), 3ul); + EXPECT_NO_FATAL_FAILURE(DPCTLKernel_Delete(Copied_KRef)); +} + TEST_P(TestDPCTLSyclKernelInterface, CheckGetWorkGroupSize) { @@ -159,7 +167,7 @@ TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxNumSubGroups) ASSERT_TRUE(axpy_mnsg != 0); } -/* +#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxSubGroupSize) { @@ -172,7 +180,7 @@ TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxSubGroupSize) ASSERT_TRUE(add_msg_sz != 0); ASSERT_TRUE(axpy_msg_sz != 0); } -*/ +#endif TEST_P(TestDPCTLSyclKernelInterface, CheckGetCompileNumSubGroups) { @@ -215,6 +223,11 @@ TEST_F(TestDPCTLSyclKernelNullArgs, CheckNumArgsNullKRef) ASSERT_EQ(DPCTLKernel_GetNumArgs(Null_KRef), -1); } +TEST_F(TestDPCTLSyclKernelNullArgs, CheckCopyNullKRef) +{ + ASSERT_TRUE(DPCTLKernel_Copy(Null_KRef) == nullptr); +} + TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetWorkGroupsSizeNullKRef) { DPCTLSyclKernelRef NullKRef = nullptr; @@ -244,14 +257,14 @@ TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetMaxNumSubGroupsNullKRef) ASSERT_EQ(DPCTLKernel_GetMaxNumSubGroups(NullKRef), 0); } -/* +#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetMaxSubGroupSizeNullKRef) { DPCTLSyclKernelRef NullKRef = nullptr; ASSERT_EQ(DPCTLKernel_GetMaxSubGroupSize(NullKRef), 0); } -*/ +#endif TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetCompileNumSubGroupsNullKRef) {