diff --git a/docs/doc_sources/api_reference/dpctl/program.rst b/docs/doc_sources/api_reference/dpctl/program.rst index 683d841541..aee4b574a8 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -20,6 +20,8 @@ execution via :py:meth:`dpctl.SyclQueue.submit`. :toctree: generated :nosignatures: + create_kernel_bundle_from_source + create_kernel_bundle_from_spirv create_program_from_source create_program_from_spirv @@ -27,11 +29,11 @@ execution via :py:meth:`dpctl.SyclQueue.submit`. :toctree: generated :nosignatures: - SyclProgram + SyclKernelBundle SyclKernel .. autosummary:: :toctree: generated :nosignatures: - SyclProgramCompilationError + SyclKernelBundleCompilationError diff --git a/docs/doc_sources/api_reference/dpctl_capi.rst b/docs/doc_sources/api_reference/dpctl_capi.rst index 5c04a8f733..ce4b435aa6 100644 --- a/docs/doc_sources/api_reference/dpctl_capi.rst +++ b/docs/doc_sources/api_reference/dpctl_capi.rst @@ -33,9 +33,9 @@ Exported typedefs .. c:struct:: PySyclKernelType -.. c:struct:: PySyclProgramObject +.. c:struct:: PySyclKernelBundleObject -.. c:struct:: PySyclProgramType +.. c:struct:: PySyclKernelBundleType To check whether a particular Python object is an instance of :py:class:`dpctl.SyclQueue`: @@ -176,19 +176,19 @@ API for :c:struct:`PySyclKernelObject` the caller remains responsible for freeing ``KRef`` as appropriate. -API for :c:struct:`PySyclProgramObject` +API for :c:struct:`PySyclKernelBundleObject` --------------------------------------- -.. c:function:: DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef(struct PySyclProgramObject *prog) +.. c:function:: DPCTLSyclKernelBundleRef SyclKernelBundle_GetKernelBundleRef(struct PySyclKernelBundleObject *prog) :param prog: Input object :returns: borrowed instance of :c:struct:`DPCTLSyclKernelBundleRef` corresponding to ``sycl::kernel_bundle`` -.. c:function:: struct PySyclProgramObject * SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef) +.. c:function:: struct PySyclKernelBundleObject * SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef) :param KBRef: instance of :c:struct:`DPCTLSyclKernelBundleRef` - :returns: new Python object of type :c:struct:`PySyclProgramType` + :returns: new Python object of type :c:struct:`PySyclKernelBundleType` Note that function does not change the ownership of the ``KBRef`` instance and the caller remains responsible for freeing ``KBRef`` as appropriate. diff --git a/docs/doc_sources/known_words.txt b/docs/doc_sources/known_words.txt index 226890759c..b68a926685 100644 --- a/docs/doc_sources/known_words.txt +++ b/docs/doc_sources/known_words.txt @@ -59,6 +59,7 @@ SyclQueue SyclContext SyclEvent SyclKernel +SyclKernelBundle SyclProgram SyclPlatform dtype diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index 6d36089bf2..577b88cd6d 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -59,7 +59,7 @@ class dpctl_capi PyTypeObject *PyMemoryUSMDeviceType_; PyTypeObject *PyMemoryUSMSharedType_; PyTypeObject *PyMemoryUSMHostType_; - PyTypeObject *PySyclProgramType_; + PyTypeObject *PySyclKernelBundleType_; PyTypeObject *PySyclKernelType_; DPCTLSyclDeviceRef (*SyclDevice_GetDeviceRef_)(PySyclDeviceObject *); @@ -89,9 +89,10 @@ class dpctl_capi DPCTLSyclKernelRef (*SyclKernel_GetKernelRef_)(PySyclKernelObject *); PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef, const char *); - DPCTLSyclKernelBundleRef (*SyclProgram_GetKernelBundleRef_)( - PySyclProgramObject *); - PySyclProgramObject *(*SyclProgram_Make_)(DPCTLSyclKernelBundleRef); + DPCTLSyclKernelBundleRef (*SyclKernelBundle_GetKernelBundleRef_)( + PySyclKernelBundleObject *); + PySyclKernelBundleObject *(*SyclKernelBundle_Make_)( + DPCTLSyclKernelBundleRef); bool PySyclDevice_Check_(PyObject *obj) const { @@ -113,9 +114,9 @@ class dpctl_capi { return PyObject_TypeCheck(obj, PySyclKernelType_) != 0; } - bool PySyclProgram_Check_(PyObject *obj) const + bool PySyclKernelBundle_Check_(PyObject *obj) const { - return PyObject_TypeCheck(obj, PySyclProgramType_) != 0; + return PyObject_TypeCheck(obj, PySyclKernelBundleType_) != 0; } ~dpctl_capi() @@ -165,7 +166,7 @@ class dpctl_capi Py_SyclQueueType_(nullptr), PySyclQueueType_(nullptr), Py_MemoryType_(nullptr), PyMemoryUSMDeviceType_(nullptr), PyMemoryUSMSharedType_(nullptr), PyMemoryUSMHostType_(nullptr), - PySyclProgramType_(nullptr), PySyclKernelType_(nullptr), + PySyclKernelBundleType_(nullptr), PySyclKernelType_(nullptr), SyclDevice_GetDeviceRef_(nullptr), SyclDevice_Make_(nullptr), SyclContext_GetContextRef_(nullptr), SyclContext_Make_(nullptr), SyclEvent_GetEventRef_(nullptr), SyclEvent_Make_(nullptr), @@ -174,8 +175,9 @@ class dpctl_capi Memory_GetContextRef_(nullptr), Memory_GetQueueRef_(nullptr), Memory_GetNumBytes_(nullptr), Memory_Make_(nullptr), SyclKernel_GetKernelRef_(nullptr), SyclKernel_Make_(nullptr), - SyclProgram_GetKernelBundleRef_(nullptr), SyclProgram_Make_(nullptr), - default_sycl_queue_{}, default_usm_memory_{}, as_usm_memory_{} + SyclKernelBundle_GetKernelBundleRef_(nullptr), + SyclKernelBundle_Make_(nullptr), default_sycl_queue_{}, + default_usm_memory_{}, as_usm_memory_{} { // Import Cython-generated C-API for dpctl @@ -198,7 +200,7 @@ class dpctl_capi this->PyMemoryUSMDeviceType_ = &PyMemoryUSMDeviceType; this->PyMemoryUSMSharedType_ = &PyMemoryUSMSharedType; this->PyMemoryUSMHostType_ = &PyMemoryUSMHostType; - this->PySyclProgramType_ = &PySyclProgramType; + this->PySyclKernelBundleType_ = &PySyclKernelBundleType; this->PySyclKernelType_ = &PySyclKernelType; // SyclDevice API @@ -228,8 +230,9 @@ class dpctl_capi // dpctl.program API this->SyclKernel_GetKernelRef_ = SyclKernel_GetKernelRef; this->SyclKernel_Make_ = SyclKernel_Make; - this->SyclProgram_GetKernelBundleRef_ = SyclProgram_GetKernelBundleRef; - this->SyclProgram_Make_ = SyclProgram_Make; + this->SyclKernelBundle_GetKernelBundleRef_ = + SyclKernelBundle_GetKernelBundleRef; + this->SyclKernelBundle_Make_ = SyclKernelBundle_Make; // create shared pointers to python objects used in type-casters // for dpctl::memory::usm_memory @@ -484,7 +487,7 @@ template <> struct type_caster /* This type caster associates * ``sycl::kernel_bundle`` C++ class with - * :class:`dpctl.program.SyclProgram` for the purposes of generation of + * :class:`dpctl.program.SyclKernelBundle` for the purposes of generation of * Python bindings by pybind11. */ template <> @@ -495,10 +498,10 @@ struct type_caster> { PyObject *source = src.ptr(); auto const &api = ::dpctl::detail::dpctl_capi::get(); - if (api.PySyclProgram_Check_(source)) { + if (api.PySyclKernelBundle_Check_(source)) { DPCTLSyclKernelBundleRef KBRef = - api.SyclProgram_GetKernelBundleRef_( - reinterpret_cast(source)); + api.SyclKernelBundle_GetKernelBundleRef_( + reinterpret_cast(source)); value = std::make_unique< sycl::kernel_bundle>( *(reinterpret_cast< @@ -508,7 +511,7 @@ struct type_caster> } else { throw py::type_error("Input is of unexpected type, expected " - "dpctl.program.SyclProgram"); + "dpctl.program.SyclKernelBundle"); } } @@ -517,13 +520,13 @@ struct type_caster> handle) { auto const &api = ::dpctl::detail::dpctl_capi::get(); - auto tmp = api.SyclProgram_Make_( + auto tmp = api.SyclKernelBundle_Make_( reinterpret_cast(&src)); return handle(reinterpret_cast(tmp)); } DPCTL_TYPE_CASTER(sycl::kernel_bundle, - _("dpctl.program.SyclProgram")); + _("dpctl.program.SyclKernelBundle")); }; /* This type caster associates diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index 4904d29a7c..9f547b46cd 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -23,16 +23,35 @@ from ._program import ( SyclKernel, - SyclProgram, - SyclProgramCompilationError, + SyclKernelBundle, + SyclKernelBundleCompilationError, + create_kernel_bundle_from_source, + create_kernel_bundle_from_spirv, create_program_from_source, create_program_from_spirv, ) __all__ = [ + "create_kernel_bundle_from_source", + "create_kernel_bundle_from_spirv", "create_program_from_source", "create_program_from_spirv", "SyclKernel", + "SyclKernelBundle", + "SyclKernelBundleCompilationError", "SyclProgram", - "SyclProgramCompilationError", ] + + +def __getattr__(name): + if name == "SyclProgram": + from warnings import warn + + warn( + "dpctl.program.SyclProgram is deprecated and will be removed in a " + "future release. Use dpctl.program.SyclKernelBundle instead.", + DeprecationWarning, + stacklevel=2, + ) + return SyclKernelBundle + raise AttributeError(f"module {__name__} has no attribute {name}") diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index dc4208a29b..435ef68521 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -40,22 +40,32 @@ cdef api class SyclKernel [object PySyclKernelObject, type PySyclKernelType]: cdef SyclKernel _create (DPCTLSyclKernelRef kref, str name) -cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]: +cdef api class SyclKernelBundle [ + object PySyclKernelBundleObject, type PySyclKernelBundleType +]: """ Wraps a sycl::kernel_bundle object created by using SYCL interoperability layer for OpenCL and Level-Zero backends. - SyclProgram exposes the C API from dpctl_sycl_kernel_bundle_interface.h. - A SyclProgram can be created from either a source string or a SPIR-V + SyclKernelBundle exposes the C API from + dpctl_sycl_kernel_bundle_interface.h. + A SyclKernelBundle can be created from either a source string or a SPIR-V binary file. """ - cdef DPCTLSyclKernelBundleRef _program_ref + cdef DPCTLSyclKernelBundleRef _kernel_bundle_ref @staticmethod - cdef SyclProgram _create (DPCTLSyclKernelBundleRef pref) - cdef DPCTLSyclKernelBundleRef get_program_ref (self) + cdef SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref) + cdef DPCTLSyclKernelBundleRef get_kernel_bundle_ref (self) cpdef SyclKernel get_sycl_kernel(self, str kernel_name) +cpdef create_kernel_bundle_from_source ( + SyclQueue q, unicode source, unicode copts=* +) +cpdef create_kernel_bundle_from_spirv ( + SyclQueue q, const unsigned char[:] IL, unicode copts=* +) cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) -cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL, - 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 3859314505..8737be4762 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -18,15 +18,18 @@ # cython: language_level=3 # cython: linetrace=True -"""Implements a Python interface for SYCL's program and kernel runtime classes. +"""Implements a Python interface for SYCL's kernel bundle and kernel runtime +classes. -The module also provides functions to create a SYCL program from either -a OpenCL source string or a SPIR-V binary file. +The module also provides functions to create a SYCL kernel bundle from either +an OpenCL source string or a SPIR-V binary file. """ from libc.stdint cimport uint32_t +import warnings + from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_Copy, DPCTLKernel_Delete, @@ -51,14 +54,14 @@ from dpctl._backend cimport ( # noqa: E211, E402; ) __all__ = [ - "create_program_from_source", - "create_program_from_spirv", + "create_kernel_bundle_from_source", + "create_kernel_bundle_from_spirv", "SyclKernel", - "SyclProgram", - "SyclProgramCompilationError", + "SyclKernelBundle", + "SyclKernelBundleCompilationError", ] -cdef class SyclProgramCompilationError(Exception): +cdef class SyclKernelBundleCompilationError(Exception): """This exception is raised when a ``sycl::kernel_bundle`` could not be built from either a SPIR-V binary file or a string source. """ @@ -185,38 +188,38 @@ cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef, const char *name): return SyclKernel._create(copied_KRef, name.decode("utf-8")) -cdef class SyclProgram: +cdef class SyclKernelBundle: """ Wraps a ``sycl::kernel_bundle`` object created using SYCL interoperability layer with underlying backends. Only the OpenCL and Level-Zero backends are currently supported. - SyclProgram exposes the C API from ``dpctl_sycl_kernel_bundle_interface.h``. - A SyclProgram can be created from either a source string or a SPIR-V - binary file. + SyclKernelBundle exposes the C API from + ``dpctl_sycl_kernel_bundle_interface.h``. A SyclKernelBundle can be + created from either a source string or a SPIR-V binary file. """ @staticmethod - cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef): - cdef SyclProgram ret = SyclProgram.__new__(SyclProgram) - ret._program_ref = KBRef + cdef SyclKernelBundle _create(DPCTLSyclKernelBundleRef KBRef): + cdef SyclKernelBundle ret = SyclKernelBundle.__new__(SyclKernelBundle) + ret._kernel_bundle_ref = KBRef return ret def __dealloc__(self): - DPCTLKernelBundle_Delete(self._program_ref) + DPCTLKernelBundle_Delete(self._kernel_bundle_ref) - cdef DPCTLSyclKernelBundleRef get_program_ref(self): - return self._program_ref + cdef DPCTLSyclKernelBundleRef get_kernel_bundle_ref(self): + return self._kernel_bundle_ref cpdef SyclKernel get_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") return SyclKernel._create( - DPCTLKernelBundle_GetKernel(self._program_ref, name), + DPCTLKernelBundle_GetKernel(self._kernel_bundle_ref, name), kernel_name ) def has_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") - return DPCTLKernelBundle_HasKernel(self._program_ref, name) + return DPCTLKernelBundle_HasKernel(self._kernel_bundle_ref, name) def addressof_ref(self): """Returns the address of the C API DPCTLSyclKernelBundleRef pointer @@ -224,14 +227,35 @@ cdef class SyclProgram: Returns: The address of the ``DPCTLSyclKernelBundleRef`` pointer used to - create this :class:`dpctl.SyclProgram` object cast to a ``size_t``. + create this :class:`dpctl.SyclKernelBundle` object cast to a + ``size_t``. """ - return int(self._program_ref) + return int(self._kernel_bundle_ref) -cpdef create_program_from_source(SyclQueue q, str src, str copts=""): +cdef api DPCTLSyclKernelBundleRef SyclKernelBundle_GetKernelBundleRef( + SyclKernelBundle kb +): + """ C-API function to access opaque kernel bundle reference from + Python object of type :class:`dpctl.program.SyclKernelBundle`. + """ + return kb.get_kernel_bundle_ref() + + +cdef api SyclKernelBundle SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef): + """ + C-API function to create :class:`dpctl.program.SyclKernelBundle` + instance from opaque ``sycl::kernel_bundle`` + reference. + """ + cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef) + return SyclKernelBundle._create(copied_KBRef) + + +cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""): """ - Creates a Sycl interoperability program from an OpenCL source string. + Creates a Sycl interoperability kernel bundle from an OpenCL source + string. We use the ``DPCTLKernelBundle_CreateFromOCLSource()`` C API function to create a ``sycl::kernel_bundle`` @@ -241,21 +265,21 @@ cpdef create_program_from_source(SyclQueue q, str src, str copts=""): Parameters: q (:class:`dpctl.SyclQueue`) The :class:`dpctl.SyclQueue` for which the - :class:`.SyclProgram` is going to be built. + :class:`.SyclKernelBundle` is going to be built. src (str) Source string for an OpenCL program. copts (str, optional) Optional compilation flags that will be used - when compiling the program. Default: ``""``. + when compiling the kernel bundle. Default: ``""``. Returns: - program (:class:`.SyclProgram`) - A :class:`.SyclProgram` object wrapping the + kernel_bundle (:class:`.SyclKernelBundle`) + A :class:`.SyclKernelBundle` object wrapping the ``sycl::kernel_bundle`` returned by the C API. Raises: - SyclProgramCompilationError + SyclKernelBundleCompilationError If a SYCL kernel bundle could not be created. """ @@ -269,15 +293,16 @@ cpdef create_program_from_source(SyclQueue q, str src, str copts=""): KBref = DPCTLKernelBundle_CreateFromOCLSource(CRef, DRef, Src, COpts) if KBref is NULL: - raise SyclProgramCompilationError() + raise SyclKernelBundleCompilationError() - return SyclProgram._create(KBref) + return SyclKernelBundle._create(KBref) -cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, - str copts=""): +cpdef create_kernel_bundle_from_spirv( + SyclQueue q, const unsigned char[:] IL, str copts="" +): """ - Creates a Sycl interoperability program from an SPIR-V binary. + Creates a Sycl interoperability kernel bundle from an SPIR-V binary. We use the :c:func:`DPCTLKernelBundle_CreateFromOCLSpirv` C API function to create a ``sycl::kernel_bundle`` @@ -286,21 +311,21 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, Parameters: q (:class:`dpctl.SyclQueue`) The :class:`dpctl.SyclQueue` for which the - :class:`.SyclProgram` is going to be built. + :class:`.SyclKernelBundle` is going to be built. IL (bytes) SPIR-V binary IL file for an OpenCL program. copts (str, optional) Optional compilation flags that will be used - when compiling the program. Default: ``""``. + when compiling the kernel bundle. Default: ``""``. Returns: - program (:class:`.SyclProgram`) - A :class:`.SyclProgram` object wrapping the + kernel_bundle (:class:`.SyclKernelBundle`) + A :class:`.SyclKernelBundle` object wrapping the ``sycl::kernel_bundle`` returned by the C API. Raises: - SyclProgramCompilationError + SyclKernelBundleCompilationError If a SYCL kernel bundle could not be created. """ @@ -315,25 +340,36 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, CRef, DRef, dIL, length, COpts ) if KBref is NULL: - raise SyclProgramCompilationError() + raise SyclKernelBundleCompilationError() - return SyclProgram._create(KBref) + return SyclKernelBundle._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`. +cpdef create_program_from_source(SyclQueue q, str src, str copts=""): + """This function is a deprecated alias for + :func:`dpctl.program.create_kernel_bundle_from_source`. + New code should use :func:`dpctl.program.create_kernel_bundle_from_source`. """ - return pro.get_program_ref() + warnings.warn( + "create_program_from_source is deprecated and will be removed in a " + "future release. Use create_kernel_bundle_from_source instead.", + DeprecationWarning, + stacklevel=2, + ) + return create_kernel_bundle_from_source(q, src, copts) -cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef): - """ - C-API function to create :class:`dpctl.program.SyclProgram` - instance from opaque ``sycl::kernel_bundle`` - reference. +cpdef create_program_from_spirv( + SyclQueue q, const unsigned char[:] IL, str copts="" +): + """This function is a deprecated alias for + :func:`dpctl.program.create_kernel_bundle_from_spirv`. + New code should use :func:`dpctl.program.create_kernel_bundle_from_spirv`. """ - cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef) - return SyclProgram._create(copied_KBRef) + warnings.warn( + "create_program_from_spirv is deprecated and will be removed in a " + "future release. Use create_kernel_bundle_from_spirv instead.", + DeprecationWarning, + stacklevel=2, + ) + return create_kernel_bundle_from_spirv(q, IL, copts) diff --git a/dpctl/tests/test_raw_kernel_arg.py b/dpctl/tests/test_raw_kernel_arg.py index 5dc0840835..9df42497ed 100644 --- a/dpctl/tests/test_raw_kernel_arg.py +++ b/dpctl/tests/test_raw_kernel_arg.py @@ -69,7 +69,7 @@ def launch_raw_arg_kernel(raw): spirv_file = get_spirv_abspath("raw-arg-kernel.spv") with open(spirv_file, "br") as spv: spv_bytes = spv.read() - prog = dpctl.program.create_program_from_spirv(q, spv_bytes) + prog = dpctl.program.create_kernel_bundle_from_spirv(q, spv_bytes) kernel = prog.get_sycl_kernel("__sycl_kernel_raw_arg_kernel") local_size = 16 global_size = local_size * 8 diff --git a/dpctl/tests/test_sycl_event.py b/dpctl/tests/test_sycl_event.py index 2985bf3af0..621c3ea59e 100644 --- a/dpctl/tests/test_sycl_event.py +++ b/dpctl/tests/test_sycl_event.py @@ -37,7 +37,7 @@ def produce_event(profiling=False): q = dpctl.SyclQueue("opencl:cpu", property="enable_profiling") else: q = dpctl.SyclQueue("opencl:cpu") - prog = dpctl_prog.create_program_from_source(q, oclSrc) + prog = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) addKernel = prog.get_sycl_kernel("add") n = 1024 * 1024 @@ -158,7 +158,7 @@ def test_get_wait_list(): size_t index = get_global_id(0); \ a[index] = sin(a[index]); \ }" - prog = dpctl_prog.create_program_from_source(q, oclSrc) + prog = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) addKernel = prog.get_sycl_kernel("add_k") sqrtKernel = prog.get_sycl_kernel("sqrt_k") sinKernel = prog.get_sycl_kernel("sin_k") diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index f58e869ebd..b2ee8b28a6 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -40,7 +40,7 @@ ("double", np.dtype("f8"), ctypes.c_double), ], ) -def test_create_program_from_source(ctype_str, dtype, ctypes_ctor): +def test_create_kernel_bundle_from_source(ctype_str, dtype, ctypes_ctor): try: q = dpctl.SyclQueue("opencl", property="enable_profiling") except dpctl.SyclQueueCreationError: @@ -59,7 +59,7 @@ def test_create_program_from_source(ctype_str, dtype, ctypes_ctor): " c[index] = d * a[index] + b[index];" "}" ) - prog = dpctl_prog.create_program_from_source(q, oclSrc) + prog = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) axpyKernel = prog.get_sycl_kernel("axpy") n_elems = 1024 * 512 @@ -174,7 +174,7 @@ def test_submit_async(): " (arg1[index] < arg2[index]) ? arg1[index] : arg2[index];" "}" ) - prog = dpctl_prog.create_program_from_source(q, oclSrc) + prog = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) kern1Kernel = prog.get_sycl_kernel("kern1") kern2Kernel = prog.get_sycl_kernel("kern2") kern3Kernel = prog.get_sycl_kernel("kern3") @@ -322,7 +322,7 @@ def test_submit_local_accessor_arg(): fn = get_spirv_abspath("local_accessor_kernel_inttys_fp32.spv") with open(fn, "br") as f: spirv_bytes = f.read() - prog = dpctl_prog.create_program_from_spirv(q, spirv_bytes) + prog = dpctl_prog.create_kernel_bundle_from_spirv(q, spirv_bytes) krn = prog.get_sycl_kernel("_ZTS14SyclKernel_SLMIlE") lws = 32 gws = lws * 10 diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 4b7102c264..55e855d739 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -14,7 +14,7 @@ # See the License for the specific language governing permissions and # limitations under the License. -"""Defines unit test cases for the SyclProgram and SyclKernel classes""" +"""Defines unit test cases for the SyclKernelBundle and SyclKernel classes""" import os @@ -30,23 +30,23 @@ def get_spirv_abspath(fn): return spirv_file -def _check_cpython_api_SyclProgram_GetKernelBundleRef(sycl_prog): +def _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(sycl_prog): """Checks Cython-generated C-API function - `SyclProgram_GetKernelBundleRef` defined in _program.pyx""" + `SyclKernelBundle_GetKernelBundleRef` defined in _program.pyx""" import ctypes import sys - assert type(sycl_prog) is dpctl_prog.SyclProgram + assert type(sycl_prog) is dpctl_prog.SyclKernelBundle 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" + # get capsule storing SyclKernelBundle_GetKernelBundleRef function ptr + kb_ref_fn_cap = mod.__pyx_capi__["SyclKernelBundle_GetKernelBundleRef"] + # construct Python callable to invoke "SyclKernelBundle_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 *)", + b"DPCTLSyclKernelBundleRef (struct PySyclKernelBundleObject *)", ) # PYFUNCTYPE(result_type, *arg_types) callable_maker = ctypes.PYFUNCTYPE(ctypes.c_void_p, ctypes.py_object) @@ -57,23 +57,23 @@ def _check_cpython_api_SyclProgram_GetKernelBundleRef(sycl_prog): assert r1 == r2 -def _check_cpython_api_SyclProgram_Make(sycl_prog): +def _check_cpython_api_SyclKernelBundle_Make(sycl_prog): """Checks Cython-generated C-API function - `SyclProgram_Make` defined in _program.pyx""" + `SyclKernelBundle_Make` defined in _program.pyx""" import ctypes import sys - assert type(sycl_prog) is dpctl_prog.SyclProgram + assert type(sycl_prog) is dpctl_prog.SyclKernelBundle 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" + # get capsule storing SyclKernelBundle_Make function ptr + make_prog_fn_cap = mod.__pyx_capi__["SyclKernelBundle_Make"] + # construct Python callable to invoke "SyclKernelBundle_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)", + b"struct PySyclKernelBundleObject *(DPCTLSyclKernelBundleRef)", ) # PYFUNCTYPE(result_type, *arg_types) callable_maker = ctypes.PYFUNCTYPE(ctypes.py_object, ctypes.c_void_p) @@ -148,7 +148,7 @@ def _check_cpython_api_SyclKernel_Make(krn): def _check_multi_kernel_program(prog): - assert type(prog) is dpctl_prog.SyclProgram + assert type(prog) is dpctl_prog.SyclKernelBundle assert type(prog.addressof_ref()) is int assert prog.has_sycl_kernel("add") @@ -185,11 +185,11 @@ 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) + _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(prog) + _check_cpython_api_SyclKernelBundle_Make(prog) -def test_create_program_from_source_ocl(): +def test_create_kernel_bundle_from_source_ocl(): oclSrc = " \ kernel void add(global int* a, global int* b, global int* c) { \ size_t index = get_global_id(0); \ @@ -203,11 +203,11 @@ def test_create_program_from_source_ocl(): q = dpctl.SyclQueue("opencl") except dpctl.SyclQueueCreationError: pytest.skip("No OpenCL queue is available") - prog = dpctl_prog.create_program_from_source(q, oclSrc) + prog = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) _check_multi_kernel_program(prog) -def test_create_program_from_spirv_ocl(): +def test_create_kernel_bundle_from_spirv_ocl(): try: q = dpctl.SyclQueue("opencl") except dpctl.SyclQueueCreationError: @@ -215,11 +215,11 @@ def test_create_program_from_spirv_ocl(): spirv_file = get_spirv_abspath("multi_kernel.spv") with open(spirv_file, "rb") as fin: spirv = fin.read() - prog = dpctl_prog.create_program_from_spirv(q, spirv) + prog = dpctl_prog.create_kernel_bundle_from_spirv(q, spirv) _check_multi_kernel_program(prog) -def test_create_program_from_spirv_l0(): +def test_create_kernel_bundle_from_spirv_l0(): try: q = dpctl.SyclQueue("level_zero") except dpctl.SyclQueueCreationError: @@ -227,14 +227,14 @@ def test_create_program_from_spirv_l0(): spirv_file = get_spirv_abspath("multi_kernel.spv") with open(spirv_file, "rb") as fin: spirv = fin.read() - prog = dpctl_prog.create_program_from_spirv(q, spirv) + prog = dpctl_prog.create_kernel_bundle_from_spirv(q, spirv) _check_multi_kernel_program(prog) @pytest.mark.xfail( reason="Level-zero backend does not support compilation from source" ) -def test_create_program_from_source_l0(): +def test_create_kernel_bundle_from_source_l0(): try: q = dpctl.SyclQueue("level_zero") except dpctl.SyclQueueCreationError: @@ -248,11 +248,11 @@ def test_create_program_from_source_l0(): size_t index = get_global_id(0); \ c[index] = a[index] + d*b[index]; \ }" - prog = dpctl_prog.create_program_from_source(q, oclSrc) + prog = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) _check_multi_kernel_program(prog) -def test_create_program_from_invalid_src_ocl(): +def test_create_kernel_bundle_from_invalid_src_ocl(): try: q = dpctl.SyclQueue("opencl") except dpctl.SyclQueueCreationError: @@ -260,5 +260,5 @@ def test_create_program_from_invalid_src_ocl(): invalid_oclSrc = " \ kernel void add( \ }" - with pytest.raises(dpctl_prog.SyclProgramCompilationError): - dpctl_prog.create_program_from_source(q, invalid_oclSrc) + with pytest.raises(dpctl_prog.SyclKernelBundleCompilationError): + dpctl_prog.create_kernel_bundle_from_source(q, invalid_oclSrc) diff --git a/dpctl/tests/test_work_group_memory.py b/dpctl/tests/test_work_group_memory.py index 17b689ee0a..5f60408843 100644 --- a/dpctl/tests/test_work_group_memory.py +++ b/dpctl/tests/test_work_group_memory.py @@ -62,7 +62,7 @@ def test_submit_work_group_memory(): spirv_file = get_spirv_abspath("work-group-memory-kernel.spv") with open(spirv_file, "br") as spv: spv_bytes = spv.read() - prog = dpctl.program.create_program_from_spirv(q, spv_bytes) + prog = dpctl.program.create_kernel_bundle_from_spirv(q, spv_bytes) kernel = prog.get_sycl_kernel("__sycl_kernel_local_mem_kernel") local_size = 16 global_size = local_size * 8 diff --git a/dpctl/tests/test_work_group_memory_opencl.py b/dpctl/tests/test_work_group_memory_opencl.py index b206ed7cab..786056b6f1 100644 --- a/dpctl/tests/test_work_group_memory_opencl.py +++ b/dpctl/tests/test_work_group_memory_opencl.py @@ -45,7 +45,7 @@ def test_submit_work_group_memory_opencl(): except dpctl.SyclQueueCreationError: pytest.skip("OpenCL queue could not be created") - prog = dpctl.program.create_program_from_source(q, ocl_kernel_src) + prog = dpctl.program.create_kernel_bundle_from_source(q, ocl_kernel_src) kernel = prog.get_sycl_kernel("local_mem_kernel") local_size = 16 global_size = local_size * 8 diff --git a/examples/pybind11/use_dpctl_sycl_kernel/example.py b/examples/pybind11/use_dpctl_sycl_kernel/example.py index f84124cfed..a164d2cd0b 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/example.py +++ b/examples/pybind11/use_dpctl_sycl_kernel/example.py @@ -31,7 +31,7 @@ il = fh.read() # Build the program for the selected device -pr = dppr.create_program_from_spirv(q, il, "") +pr = dppr.create_kernel_bundle_from_spirv(q, il, "") assert pr.has_sycl_kernel("double_it") # Retrieve the kernel from the problem 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 index e541861b84..57cfa2a23a 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py +++ b/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py @@ -45,7 +45,7 @@ def test_kernel_can_be_found(): q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: pytest.skip("Could not create default queue") - pr = dppr.create_program_from_spirv(q, il, "") + pr = dppr.create_kernel_bundle_from_spirv(q, il, "") assert pr.has_sycl_kernel("double_it") @@ -57,7 +57,7 @@ def test_kernel_submit_through_extension(): q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: pytest.skip("Could not create default queue") - pr = dppr.create_program_from_spirv(q, il, "") + pr = dppr.create_kernel_bundle_from_spirv(q, il, "") krn = pr.get_sycl_kernel("double_it") assert krn.num_args == 2 diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_platform_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_platform_interface.h index e803e11071..36fd573c63 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_platform_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_platform_interface.h @@ -89,7 +89,7 @@ __dpctl_give DPCTLSyclPlatformRef DPCTLPlatform_CreateFromSelector( __dpctl_keep const DPCTLSyclDeviceSelectorRef DSRef); /*! - * @brief Deletes the DPCTLSyclProgramRef pointer. + * @brief Deletes the DPCTLSyclPlatformRef pointer. * * @param PRef An opaque pointer to a sycl::platform. * @ingroup PlatformInterface