From 35df1717b16e5861a5e7e1aec3d772cc28fe034a Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Fri, 4 Nov 2022 11:23:07 -0700 Subject: [PATCH 01/15] C API for SyclProgram and SyclKernel classes --- dpctl/_backend.pxd | 2 + dpctl/apis/include/dpctl4pybind11.hpp | 95 +++++++++++++++++++ dpctl/apis/include/dpctl_capi.h | 5 + dpctl/program/_program.pxd | 5 +- dpctl/program/_program.pyx | 29 ++++++ .../dpctl_sycl_kernel_bundle_interface.h | 12 +++ .../include/dpctl_sycl_kernel_interface.h | 12 +++ .../dpctl_sycl_kernel_bundle_interface.cpp | 18 ++++ .../source/dpctl_sycl_kernel_interface.cpp | 18 ++++ 9 files changed, 194 insertions(+), 2 deletions(-) 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..54c323a2c5 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/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 default_usm_memory; std::shared_ptr default_usm_ndarray; std::shared_ptr as_usm_memory; + std::shared_ptr default_sycl_kernel; + std::shared_ptr 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 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 &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 &api = ::dpctl::detail::dpctl_capi::get(); + auto tmp = + api.SyclKernel_Make_(reinterpret_cast(&src)); + 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 &api = ::dpctl::detail::dpctl_capi::get(); + if (api.PySyclProgram_Check_(source)) { + DPCTLSyclKernelBundleRef KBRef = api.SyclProgram_GetKernelBundleRef_( + reinterpret_cast(source)); + value = std::make_unique>( + *(reinterpret_cast *>(KBRef))); + return true; + } + else { + throw py::type_error( + "Input is of unexpected type, expected dpctl.SyclEvent"); + } + } + + static handle cast(sycl::kernel_bundle src, return_value_policy, handle) + { + auto &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 diff --git a/dpctl/apis/include/dpctl_capi.h b/dpctl/apis/include/dpctl_capi.h index d6c104581a..ab3e008cf6 100644 --- a/dpctl/apis/include/dpctl_capi.h +++ b/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; } diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 2133a7a667..933f4bbcde 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. @@ -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=*) + diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 26a9c67a7e..d013295079 100644 --- a/dpctl/program/_program.pyx +++ b/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`` 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) \ No newline at end of file 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..c747cf5773 100644 --- a/libsyclinterface/include/dpctl_sycl_kernel_interface.h +++ b/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()`. diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index e540259bb9..03718523ac 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/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); + 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 9f5e278e3b..3fc80056e8 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/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) { From be9ea33a4d414ee39d0036d952d8fb5f389ec2ea Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 4 Nov 2022 14:09:52 -0500 Subject: [PATCH 02/15] Applied pre-commit linter --- dpctl/apis/include/dpctl4pybind11.hpp | 36 ++++++++++++------- dpctl/apis/include/dpctl_capi.h | 1 - dpctl/program/_program.pxd | 1 - dpctl/program/_program.pyx | 6 ++-- .../dpctl_sycl_kernel_bundle_interface.cpp | 8 +++-- 5 files changed, 31 insertions(+), 21 deletions(-) diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 54c323a2c5..433764346d 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -100,7 +100,8 @@ struct dpctl_capi DPCTLSyclKernelRef (*SyclKernel_GetKernelRef_)(PySyclKernelObject *); PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef); - DPCTLSyclKernelBundleRef (*SyclProgram_GetKernelBundleRef_)(PySyclProgramObject *); + DPCTLSyclKernelBundleRef (*SyclProgram_GetKernelBundleRef_)( + PySyclProgramObject *); PySyclProgramObject *(*SyclProgram_Make_)(DPCTLSyclKernelBundleRef); // tensor @@ -551,8 +552,8 @@ template <> struct type_caster return true; } else { - throw py::type_error( - "Input is of unexpected type, expected dpctl.program.SyclKernel"); + throw py::type_error("Input is of unexpected type, expected " + "dpctl.program.SyclKernel"); } } @@ -567,11 +568,13 @@ template <> struct type_caster DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.program.SyclKernel")); }; -/* This type caster associates ``sycl::kernel_bundle`` C++ class with +/* 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> +template <> +struct type_caster> { public: bool load(handle src, bool) @@ -579,10 +582,14 @@ template <> struct type_caster(source)); - value = std::make_unique>( - *(reinterpret_cast *>(KBRef))); + DPCTLSyclKernelBundleRef KBRef = + api.SyclProgram_GetKernelBundleRef_( + reinterpret_cast(source)); + value = std::make_unique< + sycl::kernel_bundle>( + *(reinterpret_cast< + sycl::kernel_bundle *>( + KBRef))); return true; } else { @@ -591,15 +598,18 @@ template <> struct type_caster src, return_value_policy, handle) + static handle cast(sycl::kernel_bundle src, + return_value_policy, + handle) { auto &api = ::dpctl::detail::dpctl_capi::get(); - auto tmp = - api.SyclProgram_Make_(reinterpret_cast(&src)); + auto tmp = api.SyclProgram_Make_( + reinterpret_cast(&src)); return handle(reinterpret_cast(tmp)); } - DPCTL_TYPE_CASTER(sycl::kernel_bundle, _("dpctl.program.SyclProgram")); + DPCTL_TYPE_CASTER(sycl::kernel_bundle, + _("dpctl.program.SyclProgram")); }; } // namespace detail } // namespace pybind11 diff --git a/dpctl/apis/include/dpctl_capi.h b/dpctl/apis/include/dpctl_capi.h index ab3e008cf6..9715e42b38 100644 --- a/dpctl/apis/include/dpctl_capi.h +++ b/dpctl/apis/include/dpctl_capi.h @@ -43,7 +43,6 @@ #include "../program/_program.h" #include "../program/_program_api.h" - // clang-format on /* diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 933f4bbcde..86c338aff7 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -59,4 +59,3 @@ cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]: cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL, unicode copts=*) - diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index d013295079..32c1afc6b7 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -30,8 +30,8 @@ from libc.stdint cimport uint32_t from dpctl._backend cimport ( # noqa: E211, E402; DPCTLCString_Delete, - DPCTLKernel_Delete, DPCTLKernel_Copy, + DPCTLKernel_Delete, DPCTLKernel_GetCompileNumSubGroups, DPCTLKernel_GetCompileSubGroupSize, DPCTLKernel_GetMaxNumSubGroups, @@ -39,10 +39,10 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_GetPreferredWorkGroupSizeMultiple, DPCTLKernel_GetPrivateMemSize, DPCTLKernel_GetWorkGroupSize, + DPCTLKernelBundle_Copy, DPCTLKernelBundle_CreateFromOCLSource, DPCTLKernelBundle_CreateFromSpirv, DPCTLKernelBundle_Delete, - DPCTLKernelBundle_Copy, DPCTLKernelBundle_GetKernel, DPCTLKernelBundle_HasKernel, DPCTLSyclContextRef, @@ -318,4 +318,4 @@ cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef): instance from opaque sycl kernel bundle reference. """ cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef) - return SyclProgram._create(copied_KBRef) \ No newline at end of file + return SyclProgram._create(copied_KBRef) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 03718523ac..b4455e0643 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -746,12 +746,14 @@ 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__); + error_handler( + "Cannot copy DPCTLSyclKernelBundleRef as input is a nullptr", + __FILE__, __func__, __LINE__); return nullptr; } try { - auto CopiedBundle = new kernel_bundle(*Bundle); + auto CopiedBundle = + new kernel_bundle(*Bundle); return wrap(CopiedBundle); } catch (std::exception const &e) { error_handler(e, __FILE__, __func__, __LINE__); From 0da4e4c50d029abd88baf3953b1b93a7ccd04c85 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Fri, 4 Nov 2022 12:26:51 -0700 Subject: [PATCH 03/15] Implemented DPCTLKernel_GetName --- dpctl/_backend.pxd | 1 + dpctl/program/_program.pyx | 5 ++++- .../include/dpctl_sycl_kernel_interface.h | 12 ++++++++++++ .../source/dpctl_sycl_kernel_interface.cpp | 15 +++++++++++++++ 4 files changed, 32 insertions(+), 1 deletion(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index f4ce41f6a2..ed2f86a777 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -262,6 +262,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 const char *DPCTLKernel_GetName(const DPCTLSyclKernelRef KRef) cdef void DPCTLKernel_Delete(DPCTLSyclKernelRef KRef) cdef DPCTLSyclKernelRef DPCTLKernel_Copy(const DPCTLSyclKernelRef KRef) cdef size_t DPCTLKernel_GetWorkGroupSize(const DPCTLSyclKernelRef KRef) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 32c1afc6b7..3390058d32 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -36,6 +36,7 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_GetCompileSubGroupSize, DPCTLKernel_GetMaxNumSubGroups, DPCTLKernel_GetNumArgs, + DPCTLKernel_GetName, DPCTLKernel_GetPreferredWorkGroupSizeMultiple, DPCTLKernel_GetPrivateMemSize, DPCTLKernel_GetWorkGroupSize, @@ -178,7 +179,9 @@ cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef): instance from opaque sycl kernel reference. """ cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef) - return SyclKernel._create(copied_KRef, "foo") + cdef const char *name = DPCTLKernel_GetName(copied_KRef) + copied_name = name.decode("UTF-8") + return SyclKernel._create(copied_KRef, copied_name) cdef class SyclProgram: """ Wraps a ``sycl::kernel_bundle`` object diff --git a/libsyclinterface/include/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_interface.h index c747cf5773..fd96d87a39 100644 --- a/libsyclinterface/include/dpctl_sycl_kernel_interface.h +++ b/libsyclinterface/include/dpctl_sycl_kernel_interface.h @@ -51,6 +51,18 @@ DPCTL_C_EXTERN_C_BEGIN DPCTL_API size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef); +/*! + * @brief Returns a C string for the function name. + * + * @param KRef DPCTLSyclKernelRef pointer to a SYCL + * interoperability kernel. + * @return A C string containing the name of the function. + * @ingroup KernelInterface + */ +DPCTL_API +__dpctl_give const char * +DPCTLKernel_GetName(__dpctl_keep const DPCTLSyclKernelRef KRef); + /*! * @brief Deletes the DPCTLSyclKernelRef after casting it to a * ``sycl::kernel``. diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index 3fc80056e8..0e7997010f 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -54,6 +54,21 @@ size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef) return static_cast(num_args); } +__dpctl_give const char * +DPCTLKernel_GetName(__dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Cannot get the name from " + "DPCTLSyclKernelRef as input is a nullptr.", + __FILE__, __func__, __LINE__); + return nullptr; + } + + auto sycl_kernel = unwrap(KRef); + auto name = sycl_kernel->get_info(); + return dpctl::helper::cstring_from_string(name); +} + void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef) { delete unwrap(KRef); From a3ce695255cc71d336c34810b0032f66da92d5c7 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 4 Nov 2022 14:35:09 -0500 Subject: [PATCH 04/15] SyclKernel_Make now takes KernelRef and const char * for the name. Nullptr is also handled, interpreting it as "default_name". Pybind11 caster provides a fixed name 'dpctl4pybind11_kernel' --- dpctl/apis/include/dpctl4pybind11.hpp | 5 +++-- dpctl/program/_program.pyx | 16 +++++++++++----- 2 files changed, 14 insertions(+), 7 deletions(-) diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 433764346d..7604964b2c 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -98,7 +98,7 @@ struct dpctl_capi // program DPCTLSyclKernelRef (*SyclKernel_GetKernelRef_)(PySyclKernelObject *); - PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef); + PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef, const char *); DPCTLSyclKernelBundleRef (*SyclProgram_GetKernelBundleRef_)( PySyclProgramObject *); @@ -561,7 +561,8 @@ template <> struct type_caster { auto &api = ::dpctl::detail::dpctl_capi::get(); auto tmp = - api.SyclKernel_Make_(reinterpret_cast(&src)); + api.SyclKernel_Make_(reinterpret_cast(&src), + "dpctl4pybind11_kernel"); return handle(reinterpret_cast(tmp)); } diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 3390058d32..5e444f7060 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -35,8 +35,8 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_GetCompileNumSubGroups, DPCTLKernel_GetCompileSubGroupSize, DPCTLKernel_GetMaxNumSubGroups, - DPCTLKernel_GetNumArgs, DPCTLKernel_GetName, + DPCTLKernel_GetNumArgs, DPCTLKernel_GetPreferredWorkGroupSizeMultiple, DPCTLKernel_GetPrivateMemSize, DPCTLKernel_GetWorkGroupSize, @@ -167,21 +167,25 @@ 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): + +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) - cdef const char *name = DPCTLKernel_GetName(copied_KRef) - copied_name = name.decode("UTF-8") - return SyclKernel._create(copied_KRef, copied_name) + 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 @@ -309,12 +313,14 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, 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` From 8c20dc483f82f760d11afebe30e062455213d782 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Fri, 4 Nov 2022 15:46:34 -0500 Subject: [PATCH 05/15] Added tests for cpython_api for SyclProgram and SyclKernel New checks are complete with docstrings, check for SyclKernel_Make also checks the nullptr for name usage. --- dpctl/tests/test_sycl_program.py | 123 +++++++++++++++++++++++++++++++ 1 file changed, 123 insertions(+) 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 = " \ From 90a453a9c332fd18eb9e4f9c7b8b8453628bc171 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sat, 5 Nov 2022 17:26:09 -0500 Subject: [PATCH 06/15] Adjusted use of wrap/unwrap to templated as per changes in master --- .../source/dpctl_sycl_kernel_bundle_interface.cpp | 4 ++-- libsyclinterface/source/dpctl_sycl_kernel_interface.cpp | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index f50c4fb299..b196264350 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -743,7 +743,7 @@ void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef) __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef) { - auto Bundle = unwrap(KBRef); + auto Bundle = unwrap>(KBRef); if (!Bundle) { error_handler( "Cannot copy DPCTLSyclKernelBundleRef as input is a nullptr", @@ -753,7 +753,7 @@ DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef) try { auto CopiedBundle = new kernel_bundle(*Bundle); - return wrap(CopiedBundle); + 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 236045c461..6520bfc70b 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -62,7 +62,7 @@ DPCTLKernel_GetName(__dpctl_keep const DPCTLSyclKernelRef KRef) return nullptr; } - auto sycl_kernel = unwrap(KRef); + auto sycl_kernel = unwrap(KRef); auto name = sycl_kernel->get_info(); return dpctl::helper::cstring_from_string(name); } @@ -75,7 +75,7 @@ void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef) __dpctl_give DPCTLSyclKernelRef DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef) { - auto Kernel = unwrap(KRef); + auto Kernel = unwrap(KRef); if (!Kernel) { error_handler("Cannot copy DPCTLSyclKernelRef as input is a nullptr", __FILE__, __func__, __LINE__); @@ -83,7 +83,7 @@ DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef) } try { auto CopiedKernel = new kernel(*Kernel); - return wrap(CopiedKernel); + return wrap(CopiedKernel); } catch (std::exception const &e) { error_handler(e, __FILE__, __func__, __LINE__); return nullptr; From 69d3c7e83baf04be039540542dd0f3de362b0f83 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sat, 5 Nov 2022 23:15:23 -0500 Subject: [PATCH 07/15] Removed DPCTLKernel_GetName as unused --- dpctl/_backend.pxd | 1 - dpctl/program/_program.pyx | 1 - .../include/dpctl_sycl_kernel_interface.h | 12 ------------ .../source/dpctl_sycl_kernel_interface.cpp | 15 --------------- 4 files changed, 29 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index ed2f86a777..f4ce41f6a2 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -262,7 +262,6 @@ 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 const char *DPCTLKernel_GetName(const DPCTLSyclKernelRef KRef) cdef void DPCTLKernel_Delete(DPCTLSyclKernelRef KRef) cdef DPCTLSyclKernelRef DPCTLKernel_Copy(const DPCTLSyclKernelRef KRef) cdef size_t DPCTLKernel_GetWorkGroupSize(const DPCTLSyclKernelRef KRef) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 5e444f7060..82ff39de56 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -35,7 +35,6 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_GetCompileNumSubGroups, DPCTLKernel_GetCompileSubGroupSize, DPCTLKernel_GetMaxNumSubGroups, - DPCTLKernel_GetName, DPCTLKernel_GetNumArgs, DPCTLKernel_GetPreferredWorkGroupSizeMultiple, DPCTLKernel_GetPrivateMemSize, diff --git a/libsyclinterface/include/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_interface.h index fd96d87a39..c747cf5773 100644 --- a/libsyclinterface/include/dpctl_sycl_kernel_interface.h +++ b/libsyclinterface/include/dpctl_sycl_kernel_interface.h @@ -51,18 +51,6 @@ DPCTL_C_EXTERN_C_BEGIN DPCTL_API size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef); -/*! - * @brief Returns a C string for the function name. - * - * @param KRef DPCTLSyclKernelRef pointer to a SYCL - * interoperability kernel. - * @return A C string containing the name of the function. - * @ingroup KernelInterface - */ -DPCTL_API -__dpctl_give const char * -DPCTLKernel_GetName(__dpctl_keep const DPCTLSyclKernelRef KRef); - /*! * @brief Deletes the DPCTLSyclKernelRef after casting it to a * ``sycl::kernel``. diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index 6520bfc70b..05db9582c0 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -52,21 +52,6 @@ size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef) return static_cast(num_args); } -__dpctl_give const char * -DPCTLKernel_GetName(__dpctl_keep const DPCTLSyclKernelRef KRef) -{ - if (!KRef) { - error_handler("Cannot get the name from " - "DPCTLSyclKernelRef as input is a nullptr.", - __FILE__, __func__, __LINE__); - return nullptr; - } - - auto sycl_kernel = unwrap(KRef); - auto name = sycl_kernel->get_info(); - return dpctl::helper::cstring_from_string(name); -} - void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef) { delete unwrap(KRef); From 35f93a3c5e772691ada7714894fc50c65a119d91 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sun, 6 Nov 2022 07:44:39 -0600 Subject: [PATCH 08/15] Make sure _program.h and _program_api.h are copied to the layout --- MANIFEST.in | 2 ++ 1 file changed, 2 insertions(+) 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 * From ea193603ae592428dae020a6ba29a88cf13b5f32 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sun, 6 Nov 2022 08:20:55 -0600 Subject: [PATCH 09/15] Added libsyclinterface/tests for DCPTLKernel_Copy and DPCTLKernelBundle_Copy Also change commented out, or preprocessor disabled code for testing, implementing, or declaring DCPTLKernel_GetMaxSubGroupsSize since it is now supported in 2023 compiler. --- .../include/dpctl_sycl_kernel_interface.h | 3 ++- .../source/dpctl_sycl_kernel_interface.cpp | 8 ++++--- .../test_sycl_kernel_bundle_interface.cpp | 21 +++++++++++++++++ .../tests/test_sycl_kernel_interface.cpp | 23 +++++++++++++++---- 4 files changed, 46 insertions(+), 9 deletions(-) diff --git a/libsyclinterface/include/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_interface.h index c747cf5773..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" @@ -129,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_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index 05db9582c0..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" @@ -160,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 @@ -179,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) { From 50268ddb93aa46fa7211c8d1a3d4d90733ba8b4e Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sun, 6 Nov 2022 09:54:17 -0600 Subject: [PATCH 10/15] Condition test exercising fp16 support on availability of HW support --- dpctl/tests/test_usm_ndarray_ctor.py | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) 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 From 753548bbcd3ec5e7a8859fca67321a16057ee140 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sun, 6 Nov 2022 15:38:00 -0600 Subject: [PATCH 11/15] Fixed text of py::type_error per PR review --- dpctl/apis/include/dpctl4pybind11.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 7604964b2c..77845883cb 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -594,8 +594,8 @@ struct type_caster> return true; } else { - throw py::type_error( - "Input is of unexpected type, expected dpctl.SyclEvent"); + throw py::type_error("Input is of unexpected type, expected " + "dpctl.program.SyclProgram"); } } From 375fa7711a2475572b578a46016252ca5357a983 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Mon, 7 Nov 2022 08:55:08 -0600 Subject: [PATCH 12/15] Added missing initialization of function pointers for dpctl.program CAPI functions dpctl_capi constructor must initialize DPCTLKernel_GetKernelRef_ and DPCTLKernelBundle_GetKernelBundleRef_ for casters to work correctly. --- dpctl/apis/include/dpctl4pybind11.hpp | 58 ++++++++++++++------------- 1 file changed, 30 insertions(+), 28 deletions(-) diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 77845883cb..da0aed8ae4 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -248,7 +248,9 @@ struct dpctl_capi 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 @@ -403,7 +405,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)); @@ -419,7 +421,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)); @@ -438,7 +440,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)); @@ -454,7 +456,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)); @@ -473,7 +475,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)); @@ -489,7 +491,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)); @@ -508,7 +510,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)); @@ -524,7 +526,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)); @@ -543,7 +545,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.PySyclKernel_Check_(source)) { DPCTLSyclKernelRef KRef = api.SyclKernel_GetKernelRef_( reinterpret_cast(source)); @@ -559,7 +561,7 @@ template <> struct type_caster static handle cast(sycl::kernel src, return_value_policy, handle) { - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); auto tmp = api.SyclKernel_Make_(reinterpret_cast(&src), "dpctl4pybind11_kernel"); @@ -581,7 +583,7 @@ 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.PySyclProgram_Check_(source)) { DPCTLSyclKernelBundleRef KBRef = api.SyclProgram_GetKernelBundleRef_( @@ -603,7 +605,7 @@ struct type_caster> return_value_policy, handle) { - auto &api = ::dpctl::detail::dpctl_capi::get(); + auto const &api = ::dpctl::detail::dpctl_capi::get(); auto tmp = api.SyclProgram_Make_( reinterpret_cast(&src)); return handle(reinterpret_cast(tmp)); @@ -650,7 +652,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; @@ -659,14 +661,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); } @@ -769,7 +771,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); } @@ -782,7 +784,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); } @@ -790,7 +792,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); } @@ -804,7 +806,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); } @@ -812,7 +814,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); @@ -829,7 +831,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); @@ -863,7 +865,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)); } @@ -872,7 +874,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); } @@ -880,7 +882,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); } @@ -888,28 +890,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_); } From a309aa20fc4cb80c4651823950b384556aec0f29 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Mon, 7 Nov 2022 13:27:17 -0600 Subject: [PATCH 13/15] Added an example to exercise pybind11 bindings for dpctl.program.SyclKernel --- .../use_dpctl_sycl_kernel/CMakeLists.txt | 36 +++++++ .../pybind11/use_dpctl_sycl_kernel/README.md | 29 ++++++ .../pybind11/use_dpctl_sycl_kernel/example.py | 41 ++++++++ .../use_dpctl_sycl_kernel/resource/README.md | 8 ++ .../resource/double_it.cl | 30 ++++++ .../resource/double_it.spv | Bin 0 -> 772 bytes .../pybind11/use_dpctl_sycl_kernel/setup.py | 26 +++++ .../tests/test_user_kernel.py | 69 +++++++++++++ .../use_kernel/__init__.py | 34 +++++++ .../use_kernel/_example.cpp | 92 ++++++++++++++++++ 10 files changed, 365 insertions(+) create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/README.md create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/example.py create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/resource/README.md create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/setup.py create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py create mode 100644 examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp 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..5eb8cbc1ef --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/README.md @@ -0,0 +1,29 @@ +# Usage of dpctl Entities in Pybind11 + +## Description + +This extension demonstrates how you can use dpctl Python types, +such as ``dpctl.SyclQueue``, 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 + +``` +(idp) [17:25:27 ansatnuc04 use_dpctl_syclqueue]$ python example.py +EU count returned by Pybind11 extension 24 +EU count computed by dpctl 24 + +Computing modular reduction using SYCL on a NumPy array +Offloaded result agrees with reference one computed by NumPy +``` 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..ea6129ec0a --- /dev/null +++ b/examples/pybind11/use_dpctl_sycl_kernel/example.py @@ -0,0 +1,41 @@ +# 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 + +q = dpctl.SyclQueue() + +with open("resource/double_it.spv", "br") as fh: + il = fh.read() + +pr = dppr.create_program_from_spirv(q, il, "") +assert pr.has_sycl_kernel("double_it") + +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.empty_like(x) + +eg.submit_custom_kernel(q, krn, x, y) + +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 0000000000000000000000000000000000000000..c193a1f02b46985a092249044ae4a822c36fe0f1 GIT binary patch literal 772 zcma)(%}T>i5QV2r+N!nHT7Nb!DuN)2;6??}g@U0Vh))pGpG6=hB~6Nb37^V`awB-Y zm|&OQFwC5pbLLL&IQ8R}xi#xr({}Cm^sH{KLVTn2P5HJ^P5xyzjUS&c=GnxgbBbJD z`7LXyi`lwtlH{Xk>W7(PdN{>rHv6av)5*;D=d&E+8}jdE+?HP!v1^<1(>O~j+!Y=NQ6wi_L-UvgX&NlU3G$-Xj{I9574`l~{Zg@K z9p%kVW&fFtx9~I{n*Abni#6!F^SjfqYAz`QRQYT9ch_$nf*&$+0&gA+}n{=cWonsZk^KN&<35Z4 literal 0 HcmV?d00001 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()); +} From 49efb613e212b7ccd346608702773688ada0f8c1 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Mon, 7 Nov 2022 18:00:34 -0600 Subject: [PATCH 14/15] Addressed PR feedback, remove unused struct member fields --- dpctl/apis/include/dpctl4pybind11.hpp | 34 +++++++++++++-------------- 1 file changed, 16 insertions(+), 18 deletions(-) diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index da0aed8ae4..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_; @@ -160,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: @@ -188,16 +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_kernel; - std::shared_ptr default_sycl_program; + 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 @@ -312,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 = @@ -334,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{}); } From c873896c55e6bd1dcfd287d8d4a6e27f65dc3328 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Mon, 7 Nov 2022 19:07:36 -0600 Subject: [PATCH 15/15] Added comments, and updated README.md --- examples/pybind11/use_dpctl_sycl_kernel/README.md | 12 ++++-------- examples/pybind11/use_dpctl_sycl_kernel/example.py | 8 +++++++- 2 files changed, 11 insertions(+), 9 deletions(-) diff --git a/examples/pybind11/use_dpctl_sycl_kernel/README.md b/examples/pybind11/use_dpctl_sycl_kernel/README.md index 5eb8cbc1ef..77aa57bf6e 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/README.md +++ b/examples/pybind11/use_dpctl_sycl_kernel/README.md @@ -3,8 +3,8 @@ ## Description This extension demonstrates how you can use dpctl Python types, -such as ``dpctl.SyclQueue``, in Pybind11 -extensions. +such as ``dpctl.SyclQueue`` and ``dpctl.program.SyclKernel``, in +Pybind11 extensions. ## Building @@ -20,10 +20,6 @@ python example.py # Sample output ``` -(idp) [17:25:27 ansatnuc04 use_dpctl_syclqueue]$ python example.py -EU count returned by Pybind11 extension 24 -EU count computed by dpctl 24 - -Computing modular reduction using SYCL on a NumPy array -Offloaded result agrees with reference one computed by NumPy +(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 index ea6129ec0a..292cdf414d 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/example.py +++ b/examples/pybind11/use_dpctl_sycl_kernel/example.py @@ -22,20 +22,26 @@ 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, x, y) +eg.submit_custom_kernel(q, krn, src=x, dst=y) +# output the result print(dpt.asnumpy(y))