diff --git a/.flake8 b/.flake8 index 1453996659..4d0b83f2e0 100644 --- a/.flake8 +++ b/.flake8 @@ -22,7 +22,7 @@ per-file-ignores = dpctl/_sycl_queue.pyx: E999, E225, E226, E227 dpctl/_sycl_queue_manager.pyx: E999, E225 dpctl/memory/_memory.pyx: E999, E225, E226, E227 - dpctl/program/_program.pyx: E999, E225, E226, E227 + dpctl/compiler/_compiler.pyx: E999, E225, E226, E227 dpctl/tests/_cython_api.pyx: E999, E225, E227, E402 dpctl/utils/_onetrace_context.py: E501, W505 examples/cython/sycl_buffer/syclbuffer/_syclbuffer.pyx: E999, E225, E402 diff --git a/CHANGELOG.md b/CHANGELOG.md index a183f475d1..0db1e68148 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,12 +9,18 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Added * `dpctl.SyclQueue.copy` and `dpctl.SyclQueue.copy_async` methods [gh-2273](https://github.com/IntelPython/dpctl/pull/2273) +### Removed +* Removed Cython API for `dpctl.program` submodule, as the submodule is now deprecated, with functionality migrated to `dpctl.compiler` [gh-2317](https://github.com/IntelPython/dpctl/pull/2317) + ### Changed * Bump minimum NumPy version to 1.26 [gh-2192](https://github.com/IntelPython/dpctl/pull/2192) * Rewrote USM Python examples into a single example [gh-2292](https://github.com/IntelPython/dpctl/pull/2292) ### Fixed +### Deprecated +* Deprecated `dpctl.program` submodule in favor of `dpctl.compiler`, which provides a better description of the purpose of the submodule in exposing DPC++ compilation-related functionality [gh-2317](https://github.com/IntelPython/dpctl/pull/2317) + ## [0.22.1] - Apr. 24, 2026 This is a bug-fix release which fixes a memory leak in `dpctl.RawKernelArg`. diff --git a/docs/doc_sources/api_reference/dpctl/compiler.rst b/docs/doc_sources/api_reference/dpctl/compiler.rst new file mode 100644 index 0000000000..fae4c27373 --- /dev/null +++ b/docs/doc_sources/api_reference/dpctl/compiler.rst @@ -0,0 +1,38 @@ +.. _dpctl_compiler_pyapi: + +:py:mod:`dpctl.compiler` +======================== + +:py:mod:`dpctl.compiler` provides a way to create a SYCL kernel +from either an OpenCL* program source code represented as a string +or a SPIR-V binary file. + +It implements creation of interoperability +``sycl::kernel_bundle`` (a collection of kernels), +as well as creation of individual ``sycl::kernel``, suitable for submission for +execution via :py:meth:`dpctl.SyclQueue.submit`. + +.. py:module:: dpctl.compiler + +.. currentmodule:: dpctl.compiler + +.. autosummary:: + :toctree: generated + :nosignatures: + + create_kernel_bundle_from_source + create_kernel_bundle_from_spirv + +.. autosummary:: + :toctree: generated + :nosignatures: + + SyclKernelBundle + SyclKernel + SpecializationConstant + +.. autosummary:: + :toctree: generated + :nosignatures: + + SyclKernelBundleCompilationError diff --git a/docs/doc_sources/api_reference/dpctl/index.rst b/docs/doc_sources/api_reference/dpctl/index.rst index 6af99e288f..aa84cb12b9 100644 --- a/docs/doc_sources/api_reference/dpctl/index.rst +++ b/docs/doc_sources/api_reference/dpctl/index.rst @@ -14,8 +14,10 @@ * - :py:mod:`dpctl.memory` - Unified Shared Memory operations - * - :py:mod:`dpctl.program` + * - :py:mod:`dpctl.compiler` - Support for working with SYCL kernels + * - :py:mod:`dpctl.program` + - (deprecated, use :py:mod:`dpctl.compiler`) * - :py:mod:`dpctl.utils` - A collection of utility functions diff --git a/docs/doc_sources/api_reference/dpctl/program.rst b/docs/doc_sources/api_reference/dpctl/program.rst index 6b2bf22bad..61dc1839d4 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -1,14 +1,17 @@ .. _dpctl_program_pyapi: -:py:mod:`dpctl.program` -======================= +:py:mod:`dpctl.program` (deprecated) +===================================== + +.. deprecated:: + :py:mod:`dpctl.program` is deprecated. Use :py:mod:`dpctl.compiler` instead. :py:mod:`dpctl.program` provides a way to create a SYCL kernel from either an OpenCL* program source code represented as a string or a SPIR-V binary file. It implements creation of interoperability -``sycl::kernel_bundle`` (a collection of kernels), +``sycl::kernel_bundle`` (a collection of kernels), as well as creation of individual ``sycl::kernel``, suitable for submission for execution via :py:meth:`dpctl.SyclQueue.submit`. @@ -20,33 +23,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 - -.. autosummary:: - :toctree: generated - :nosignatures: - - SyclKernelBundle SyclKernel - SpecializationConstant - -.. autosummary:: - :toctree: generated - :nosignatures: - - SyclKernelBundleCompilationError - -:py:mod:`dpctl.program.utils` ------------------------------ - -.. py:module:: dpctl.program.utils - -.. currentmodule:: dpctl.program.utils - -.. autofunction:: parse_spirv_specializations - -.. autoclass:: SpecializationConstantInfo - :members: + SyclProgram + SyclProgramCompilationError diff --git a/docs/doc_sources/api_reference/dpctl_cython.rst b/docs/doc_sources/api_reference/dpctl_cython.rst index b419bbdaff..872608167f 100644 --- a/docs/doc_sources/api_reference/dpctl_cython.rst +++ b/docs/doc_sources/api_reference/dpctl_cython.rst @@ -6,8 +6,8 @@ .. role:: python(code) :language: python -All Python modules of :py:mod:`dpctl` come with ``__init__.pxd`` alongside ``__init__.py`` files -permitting doing both :python:`import dpctl` and :code:`cimport dpctl as c_dpctl`. +Python modules of :py:mod:`dpctl` that expose a Cython API come with ``__init__.pxd`` alongside +``__init__.py`` files, permitting doing both :python:`import dpctl` and :code:`cimport dpctl as c_dpctl`. Locations of Cython declaration files in the package installation layout are as follows: @@ -24,12 +24,12 @@ Locations of Cython declaration files in the package installation layout are as _sycl_queue_manager.pxd sycl.pxd + compiler/__init__.pxd + compiler/_compiler.pxd + memory/__init__.pxd memory/_memory.pxd - program/__init__.pxd - program/_program.pxd - File ``_backend.pxd`` redefines symbols from :ref:`DPCTLSyclInterface library ` for Cython. File ``sycl.pxd`` provides casters from opaque types in "DPCTLSyclInterface" C library to SYCL C++ object pointers. diff --git a/docs/doc_sources/api_reference/index.rst b/docs/doc_sources/api_reference/index.rst index 8e301febb6..fbe3b1da63 100644 --- a/docs/doc_sources/api_reference/index.rst +++ b/docs/doc_sources/api_reference/index.rst @@ -9,7 +9,8 @@ The package ``dpctl`` provides * Python language bindings for the DPC++ runtime - :ref:`API objects ` in :py:mod:`dpctl` namespace - :ref:`API objects ` in :py:mod:`dpctl.memory` namespace - - :ref:`API objects ` in :py:mod:`dpctl.program` namespace + - :ref:`API objects ` in :py:mod:`dpctl.compiler` namespace + - :ref:`API objects ` in :py:mod:`dpctl.program` namespace (deprecated, use :py:mod:`dpctl.compiler`) - :ref:`API objects ` in :py:mod:`dpctl.utils` namespace * Python C API - :ref:`C API ` for working with Python classes defined in :mod:`dpctl` @@ -31,6 +32,7 @@ The package ``dpctl`` provides dpctl/index dpctl/memory + dpctl/compiler dpctl/program dpctl/utils libsyclinterface/index diff --git a/dpctl/CMakeLists.txt b/dpctl/CMakeLists.txt index a24c7443f9..a4378a03d0 100644 --- a/dpctl/CMakeLists.txt +++ b/dpctl/CMakeLists.txt @@ -203,6 +203,6 @@ build_dpctl_ext(${_trgt} ${_cy_file} "dpctl" SYCL) target_include_directories(${_trgt} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}) target_link_libraries(DpctlCAPI INTERFACE ${_trgt}_headers) -add_subdirectory(program) +add_subdirectory(compiler) add_subdirectory(memory) add_subdirectory(utils) diff --git a/dpctl/__init__.py b/dpctl/__init__.py index 5f57cdf90e..04391a902c 100644 --- a/dpctl/__init__.py +++ b/dpctl/__init__.py @@ -125,6 +125,7 @@ # add submodules __all__ += [ "memory", + "compiler", "program", "utils", ] diff --git a/dpctl/_sycl_queue.pxd b/dpctl/_sycl_queue.pxd index 9ca1d7c703..1e39e59279 100644 --- a/dpctl/_sycl_queue.pxd +++ b/dpctl/_sycl_queue.pxd @@ -32,7 +32,7 @@ from ._backend cimport ( from ._sycl_context cimport SyclContext from ._sycl_device cimport SyclDevice from ._sycl_event cimport SyclEvent -from .program._program cimport SyclKernel +from .compiler._compiler cimport SyclKernel cdef public api class _SyclQueue [ diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index a22586b71e..83a7ecf57d 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -386,7 +386,7 @@ kernel_arg_type = _kernel_arg_type() cdef class SyclKernelSubmitError(Exception): """ A ``SyclKernelSubmitError`` exception is raised when - the provided :class:`.program.SyclKernel` could not be + the provided :class:`.compiler.SyclKernel` could not be submitted to the :class:`.SyclQueue`. """ @@ -1249,10 +1249,10 @@ cdef class SyclQueue(_SyclQueue): list dEvents=None ): """ - Asynchronously submit :class:`dpctl.program.SyclKernel` for execution. + Asynchronously submit :class:`dpctl.compiler.SyclKernel` for execution. Args: - kernel (dpctl.program.SyclKernel): + kernel (dpctl.compiler.SyclKernel): SYCL kernel object args (List[object]): List of kernel arguments @@ -1410,10 +1410,10 @@ cdef class SyclQueue(_SyclQueue): list dEvents=None ): """ - Submit :class:`dpctl.program.SyclKernel` for execution. + Submit :class:`dpctl.compiler.SyclKernel` for execution. Args: - kernel (dpctl.program.SyclKernel): + kernel (dpctl.compiler.SyclKernel): SYCL kernel object args (List[object]): List of kernel arguments diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index e8d48694f5..f23f090de8 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -259,7 +259,7 @@ class dpctl_capi this->Memory_GetNumBytes_ = Memory_GetNumBytes; this->Memory_Make_ = Memory_Make; - // dpctl.program API + // dpctl.compiler API this->SyclKernel_GetKernelRef_ = SyclKernel_GetKernelRef; this->SyclKernel_Make_ = SyclKernel_Make; this->SyclKernelBundle_GetKernelBundleRef_ = @@ -482,7 +482,7 @@ template <> struct type_caster }; /* This type caster associates ``sycl::kernel`` C++ class with - * :class:`dpctl.program.SyclKernel` for the purposes of generation of + * :class:`dpctl.compiler.SyclKernel` for the purposes of generation of * Python bindings by pybind11. */ template <> struct type_caster @@ -501,7 +501,7 @@ template <> struct type_caster } else { throw py::type_error("Input is of unexpected type, expected " - "dpctl.program.SyclKernel"); + "dpctl.compiler.SyclKernel"); } } @@ -514,12 +514,12 @@ template <> struct type_caster return handle(reinterpret_cast(tmp)); } - DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.program.SyclKernel")); + DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.compiler.SyclKernel")); }; /* This type caster associates * ``sycl::kernel_bundle`` C++ class with - * :class:`dpctl.program.SyclKernelBundle` for the purposes of generation of + * :class:`dpctl.compiler.SyclKernelBundle` for the purposes of generation of * Python bindings by pybind11. */ template <> @@ -543,7 +543,7 @@ struct type_caster> } else { throw py::type_error("Input is of unexpected type, expected " - "dpctl.program.SyclKernelBundle"); + "dpctl.compiler.SyclKernelBundle"); } } @@ -558,7 +558,7 @@ struct type_caster> } DPCTL_TYPE_CASTER(sycl::kernel_bundle, - _("dpctl.program.SyclKernelBundle")); + _("dpctl.compiler.SyclKernelBundle")); }; /* This type caster associates diff --git a/dpctl/apis/include/dpctl_capi.h b/dpctl/apis/include/dpctl_capi.h index 9dc904a109..36e764ed49 100644 --- a/dpctl/apis/include/dpctl_capi.h +++ b/dpctl/apis/include/dpctl_capi.h @@ -45,8 +45,8 @@ #include "dpctl/_sycl_queue_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/compiler/_compiler.h" +#include "dpctl/compiler/_compiler_api.h" // clang-format on @@ -66,6 +66,6 @@ static inline void import_dpctl(void) import_dpctl___sycl_event(); import_dpctl___sycl_queue(); import_dpctl__memory___memory(); - import_dpctl__program___program(); + import_dpctl__compiler___compiler(); return; } diff --git a/dpctl/program/CMakeLists.txt b/dpctl/compiler/CMakeLists.txt similarity index 72% rename from dpctl/program/CMakeLists.txt rename to dpctl/compiler/CMakeLists.txt index e55d5b0e08..da92df546d 100644 --- a/dpctl/program/CMakeLists.txt +++ b/dpctl/compiler/CMakeLists.txt @@ -1,7 +1,6 @@ - file(GLOB _cython_sources *.pyx) foreach(_cy_file ${_cython_sources}) get_filename_component(_trgt ${_cy_file} NAME_WLE) - build_dpctl_ext(${_trgt} ${_cy_file} "dpctl/program" RELATIVE_PATH "..") + build_dpctl_ext(${_trgt} ${_cy_file} "dpctl/compiler" RELATIVE_PATH "..") target_link_libraries(DpctlCAPI INTERFACE ${_trgt}_headers) endforeach() diff --git a/dpctl/program/__init__.pxd b/dpctl/compiler/__init__.pxd similarity index 75% rename from dpctl/program/__init__.pxd rename to dpctl/compiler/__init__.pxd index 38b71beebf..ce78a4af38 100644 --- a/dpctl/program/__init__.pxd +++ b/dpctl/compiler/__init__.pxd @@ -1,6 +1,6 @@ # Data Parallel Control (dpctl) # -# Copyright 2020 Intel Corporation +# Copyright 2020-2025 Intel Corporation # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -18,8 +18,13 @@ # cython: language_level=3 """Declares the extension types and functions for the Cython API -implemented in dpctl.program._program.pyx. +implemented in dpctl.compiler._compiler.pyx. """ -from dpctl.program._program cimport * +from dpctl.compiler._compiler cimport ( + SyclKernel, + SyclKernelBundle, + create_kernel_bundle_from_source, + create_kernel_bundle_from_spirv, +) diff --git a/dpctl/compiler/__init__.py b/dpctl/compiler/__init__.py new file mode 100644 index 0000000000..3c16395140 --- /dev/null +++ b/dpctl/compiler/__init__.py @@ -0,0 +1,45 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2025 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. + +""" +**Data Parallel Control Compiler** provides a way to create a SYCL kernel +from either an OpenCL program represented as a string or a SPIR-V binary +file. + +""" + +from dpctl.compiler._compiler import ( + SpecializationConstant, + SyclKernel, + SyclKernelBundle, + SyclKernelBundleCompilationError, + create_kernel_bundle_from_source, + create_kernel_bundle_from_spirv, +) + +__all__ = [ + "create_kernel_bundle_from_source", + "create_kernel_bundle_from_spirv", + "SyclKernel", + "SyclKernelBundle", + "SyclKernelBundleCompilationError", + "SpecializationConstant", +] + +# add submodules +__all__ += [ + "utils", +] diff --git a/dpctl/program/_program.pxd b/dpctl/compiler/_compiler.pxd similarity index 90% rename from dpctl/program/_program.pxd rename to dpctl/compiler/_compiler.pxd index de75753495..47af0a604a 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/compiler/_compiler.pxd @@ -18,7 +18,7 @@ # cython: language_level=3 -"""The Cython declarations for the functions defined in _program.pyx. +"""The Cython declarations for the functions defined in _compiler.pyx. """ @@ -68,7 +68,3 @@ cpdef create_kernel_bundle_from_spirv ( unicode copts=*, list specializations=*, ) -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/compiler/_compiler.pyx similarity index 93% rename from dpctl/program/_program.pyx rename to dpctl/compiler/_compiler.pyx index 1c9227ec02..623450fd24 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/compiler/_compiler.pyx @@ -40,8 +40,6 @@ from libc.stdint cimport UINT32_MAX, uint32_t from libc.stdlib cimport free, malloc from libc.string cimport memcmp, memcpy -import warnings - from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_Copy, DPCTLKernel_Delete, @@ -189,14 +187,14 @@ cdef class SyclKernel: cdef api DPCTLSyclKernelRef SyclKernel_GetKernelRef(SyclKernel ker): """ C-API function to access opaque kernel reference from - Python object of type :class:`dpctl.program.SyclKernel`. + Python object of type :class:`dpctl.compiler.SyclKernel`. """ return ker.get_kernel_ref() cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef, const char *name): """ - C-API function to create :class:`dpctl.program.SyclKernel` + C-API function to create :class:`dpctl.compiler.SyclKernel` instance from opaque sycl kernel reference. """ cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef) @@ -255,14 +253,14 @@ 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`. + Python object of type :class:`dpctl.compiler.SyclKernelBundle`. """ return kb.get_kernel_bundle_ref() cdef api SyclKernelBundle SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef): """ - C-API function to create :class:`dpctl.program.SyclKernelBundle` + C-API function to create :class:`dpctl.compiler.SyclKernelBundle` instance from opaque ``sycl::kernel_bundle`` reference. """ @@ -275,7 +273,7 @@ cdef class SpecializationConstant: SpecializationConstant(spec_id, *args) Python class representing SYCL specialization constants that can be used - when creating a :class:`dpctl.program.SyclKernelBundle` from SPIR-V. + when creating a :class:`dpctl.compiler.SyclKernelBundle` from SPIR-V. There are multiple ways to create a :class:`.SpecializationConstant`: @@ -605,33 +603,3 @@ cpdef create_kernel_bundle_from_spirv( free(spconsts) return SyclKernelBundle._create(KBref) - - -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`. - """ - 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) - - -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`. - """ - 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/program/utils/__init__.py b/dpctl/compiler/utils/__init__.py similarity index 76% rename from dpctl/program/utils/__init__.py rename to dpctl/compiler/utils/__init__.py index 34c480838d..1359230360 100644 --- a/dpctl/program/utils/__init__.py +++ b/dpctl/compiler/utils/__init__.py @@ -1,6 +1,6 @@ # Data Parallel Control (dpctl) # -# Copyright 2026 Intel Corporation +# Copyright 2020-2025 Intel Corporation # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -15,12 +15,11 @@ # limitations under the License. """ -A collection of utility functions for dpctl.program module. +A collection of utility functions for dpctl.compiler module. """ -from ._utils import SpecializationConstantInfo, parse_spirv_specializations +from dpctl.compiler.utils._utils import parse_spirv_specializations __all__ = [ "parse_spirv_specializations", - "SpecializationConstantInfo", ] diff --git a/dpctl/program/utils/_utils.py b/dpctl/compiler/utils/_utils.py similarity index 99% rename from dpctl/program/utils/_utils.py rename to dpctl/compiler/utils/_utils.py index fe69f0448e..e1c3a6cdd3 100644 --- a/dpctl/program/utils/_utils.py +++ b/dpctl/compiler/utils/_utils.py @@ -14,7 +14,7 @@ # See the License for the specific language governing permissions and # limitations under the License. -"""Implements various utilities for the dpctl.program module.""" +"""Implements various utilities for the dpctl.compiler module.""" from dataclasses import dataclass from enum import IntEnum diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index 4eedbee7a8..bfd6c2aa1d 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -15,61 +15,58 @@ # limitations under the License. """ -**Data Parallel Control Program** provides a way to create a SYCL kernel -from either an OpenCL program represented as a string or a SPIR-V binary +**Data Parallel Control Program** (deprecated) provides a way to create a SYCL +kernel from either an OpenCL program represented as a string or a SPIR-V binary file. +.. deprecated:: + The dpctl.program module is deprecated. Use dpctl.compiler instead. + """ -from . import utils -from ._program import ( - SpecializationConstant, +import warnings + +from dpctl.compiler import ( SyclKernel, SyclKernelBundle, SyclKernelBundleCompilationError, - create_kernel_bundle_from_source, - create_kernel_bundle_from_spirv, +) +from dpctl.program._program import ( 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", - "SpecializationConstant", ] -# add submodules -__all__ += [ - "utils", -] + +warnings.warn( + "dpctl.program is deprecated and will be removed in a future release. " + "Use dpctl.compiler instead.", + DeprecationWarning, + stacklevel=2, +) def __getattr__(name): if name == "SyclProgram": - from warnings import warn - - warn( + warnings.warn( "dpctl.program.SyclProgram is deprecated and will be removed in a " - "future release. Use dpctl.program.SyclKernelBundle instead.", + "future release. Use dpctl.compiler.SyclKernelBundle instead.", DeprecationWarning, stacklevel=2, ) return SyclKernelBundle if name == "SyclProgramCompilationError": - from warnings import warn - - warn( + warnings.warn( "dpctl.program.SyclProgramCompilationError is deprecated and will " "be removed in a future release. Use " - "dpctl.program.SyclKernelBundleCompilationError instead.", + "dpctl.compiler.SyclKernelBundleCompilationError instead.", DeprecationWarning, stacklevel=2, ) diff --git a/dpctl/program/_program.py b/dpctl/program/_program.py new file mode 100644 index 0000000000..8d524ed1f5 --- /dev/null +++ b/dpctl/program/_program.py @@ -0,0 +1,60 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2025 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. + +"""Deprecated wrapper functions for backward compatibility.""" + +import warnings + +from dpctl.compiler import ( + SyclKernel, + create_kernel_bundle_from_source, + create_kernel_bundle_from_spirv, +) + + +def create_program_from_source(q, src, copts=""): + """This function is a deprecated alias for + :func:`dpctl.compiler.create_kernel_bundle_from_source`. + New code should use :func:`dpctl.compiler.create_kernel_bundle_from_source`. + """ + 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) + + +def create_program_from_spirv(q, IL, copts=""): + """This function is a deprecated alias for + :func:`dpctl.compiler.create_kernel_bundle_from_spirv`. + New code should use :func:`dpctl.compiler.create_kernel_bundle_from_spirv`. + """ + 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) + + +__all__ = [ + "create_program_from_source", + "create_program_from_spirv", + "SyclKernel", +] diff --git a/dpctl/tests/test_raw_kernel_arg.py b/dpctl/tests/test_raw_kernel_arg.py index fb28be6a30..fd07930935 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() - kb = dpctl.program.create_kernel_bundle_from_spirv(q, spv_bytes) + kb = dpctl.compiler.create_kernel_bundle_from_spirv(q, spv_bytes) kernel = kb.get_sycl_kernel("__sycl_kernel_raw_arg_kernel") local_size = 16 global_size = local_size * 8 diff --git a/dpctl/tests/test_sycl_context.py b/dpctl/tests/test_sycl_context.py index fc3dcfb52d..71deb8794a 100644 --- a/dpctl/tests/test_sycl_context.py +++ b/dpctl/tests/test_sycl_context.py @@ -168,22 +168,22 @@ def test_context_multi_device(): assert type(repr(ctx)) is str q1 = dpctl.SyclQueue(ctx, d1) q2 = dpctl.SyclQueue(ctx, d2) - import dpctl.memory as dpmem + import dpctl.memory as dpm - shmem_1 = dpmem.MemoryUSMShared(256, queue=q1) - shmem_2 = dpmem.MemoryUSMDevice(256, queue=q2) + shmem_1 = dpm.MemoryUSMShared(256, queue=q1) + shmem_2 = dpm.MemoryUSMDevice(256, queue=q2) shmem_2.copy_from_device(shmem_1) # create context for single sub-device ctx1 = dpctl.SyclContext(d1) q1 = dpctl.SyclQueue(ctx1, d1) - shmem_1 = dpmem.MemoryUSMShared(256, queue=q1) + shmem_1 = dpm.MemoryUSMShared(256, queue=q1) cap = ctx1._get_capsule() cap2 = ctx1._get_capsule() del ctx1 del cap2 # exercise deleter of non-renamed capsule ctx2 = dpctl.SyclContext(cap) q2 = dpctl.SyclQueue(ctx2, d1) - shmem_2 = dpmem.MemoryUSMDevice(256, queue=q2) + shmem_2 = dpm.MemoryUSMDevice(256, queue=q2) shmem_2.copy_from_device(shmem_1) diff --git a/dpctl/tests/test_sycl_event.py b/dpctl/tests/test_sycl_event.py index e46374d46c..01c61ca7f6 100644 --- a/dpctl/tests/test_sycl_event.py +++ b/dpctl/tests/test_sycl_event.py @@ -20,8 +20,8 @@ import pytest import dpctl -import dpctl.memory as dpctl_mem -import dpctl.program as dpctl_prog +import dpctl.compiler as dpc +import dpctl.memory as dpm from dpctl import event_status_type as esty from .helper import create_invalid_capsule @@ -37,12 +37,12 @@ def produce_event(profiling=False): q = dpctl.SyclQueue("opencl:cpu", property="enable_profiling") else: q = dpctl.SyclQueue("opencl:cpu") - kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) addKernel = kb.get_sycl_kernel("add") n = 1024 * 1024 a = np.arange(n, dtype="i") - a_usm = dpctl_mem.MemoryUSMDevice(a.nbytes, queue=q) + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) ev1 = q.memcpy_async(dest=a_usm, src=a, count=a.nbytes) args = [a_usm] @@ -158,14 +158,14 @@ def test_get_wait_list(): size_t index = get_global_id(0); \ a[index] = sin(a[index]); \ }" - kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) addKernel = kb.get_sycl_kernel("add_k") sqrtKernel = kb.get_sycl_kernel("sqrt_k") sinKernel = kb.get_sycl_kernel("sin_k") n = 1024 * 1024 a = np.arange(n, dtype="f") - a_usm = dpctl_mem.MemoryUSMDevice(a.nbytes, queue=q) + a_usm = dpm.MemoryUSMDevice(a.nbytes, queue=q) ev_1 = q.memcpy_async(dest=a_usm, src=a, count=a.nbytes) args = [a_usm] @@ -200,8 +200,8 @@ def test_sycl_timer(): except dpctl.SyclQueueCreationError: pytest.skip("Queue creation of default device failed") timer = dpctl.SyclTimer() - m1 = dpctl_mem.MemoryUSMDevice(1024 * 1024, queue=q) - m2 = dpctl_mem.MemoryUSMDevice(1024 * 1024, queue=q) + m1 = dpm.MemoryUSMDevice(1024 * 1024, queue=q) + m2 = dpm.MemoryUSMDevice(1024 * 1024, queue=q) with timer(q): # device task m1.copy_from_device(m2) diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index e2d0971111..654b988a19 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -23,8 +23,8 @@ import pytest import dpctl +import dpctl.compiler as dpc import dpctl.memory as dpm -import dpctl.program as dpctl_prog from dpctl._sycl_queue import kernel_arg_type @@ -59,7 +59,7 @@ def test_create_kernel_bundle_from_source(ctype_str, dtype, ctypes_ctor): " c[index] = d * a[index] + b[index];" "}" ) - kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) axpyKernel = kb.get_sycl_kernel("axpy") n_elems = 1024 * 512 @@ -174,14 +174,14 @@ def test_submit_async(): " (arg1[index] < arg2[index]) ? arg1[index] : arg2[index];" "}" ) - kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) kern1Kernel = kb.get_sycl_kernel("kern1") kern2Kernel = kb.get_sycl_kernel("kern2") kern3Kernel = kb.get_sycl_kernel("kern3") - assert isinstance(kern1Kernel, dpctl_prog.SyclKernel) - assert isinstance(kern2Kernel, dpctl_prog.SyclKernel) - assert isinstance(kern2Kernel, dpctl_prog.SyclKernel) + assert isinstance(kern1Kernel, dpc.SyclKernel) + assert isinstance(kern2Kernel, dpc.SyclKernel) + assert isinstance(kern2Kernel, dpc.SyclKernel) status_complete = dpctl.event_status_type.complete @@ -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() - kb = dpctl_prog.create_kernel_bundle_from_spirv(q, spirv_bytes) + kb = dpc.create_kernel_bundle_from_spirv(q, spirv_bytes) krn = kb.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 bbf16df98c..cc7292a021 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -22,8 +22,8 @@ import pytest import dpctl -import dpctl.program as dpctl_prog -from dpctl.program.utils import parse_spirv_specializations +import dpctl.compiler as dpc +from dpctl.compiler.utils import parse_spirv_specializations def get_spirv_abspath(fn): @@ -38,7 +38,7 @@ def _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(sycl_prog): import ctypes import sys - assert type(sycl_prog) is dpctl_prog.SyclKernelBundle + assert type(sycl_prog) is dpc.SyclKernelBundle mod = sys.modules[sycl_prog.__class__.__module__] # get capsule storing SyclKernelBundle_GetKernelBundleRef function ptr kb_ref_fn_cap = mod.__pyx_capi__["SyclKernelBundle_GetKernelBundleRef"] @@ -65,7 +65,7 @@ def _check_cpython_api_SyclKernelBundle_Make(sycl_prog): import ctypes import sys - assert type(sycl_prog) is dpctl_prog.SyclKernelBundle + assert type(sycl_prog) is dpc.SyclKernelBundle mod = sys.modules[sycl_prog.__class__.__module__] # get capsule storing SyclKernelBundle_Make function ptr make_prog_fn_cap = mod.__pyx_capi__["SyclKernelBundle_Make"] @@ -92,7 +92,7 @@ def _check_cpython_api_SyclKernel_GetKernelRef(krn): import ctypes import sys - assert type(krn) is dpctl_prog.SyclKernel + assert type(krn) is dpc.SyclKernel mod = sys.modules[krn.__class__.__module__] # get capsule storing SyclKernel_GetKernelRef function ptr k_ref_fn_cap = mod.__pyx_capi__["SyclKernel_GetKernelRef"] @@ -118,7 +118,7 @@ def _check_cpython_api_SyclKernel_Make(krn): import ctypes import sys - assert type(krn) is dpctl_prog.SyclKernel + assert type(krn) is dpc.SyclKernel mod = sys.modules[krn.__class__.__module__] # get capsule storing SyclKernel_Make function ptr k_make_fn_cap = mod.__pyx_capi__["SyclKernel_Make"] @@ -150,7 +150,7 @@ def _check_cpython_api_SyclKernel_Make(krn): def _check_multi_kernel_program(kb): - assert type(kb) is dpctl_prog.SyclKernelBundle + assert type(kb) is dpc.SyclKernelBundle assert type(kb.addressof_ref()) is int assert kb.has_sycl_kernel("add") @@ -205,7 +205,7 @@ def test_create_kernel_bundle_from_source_ocl(): q = dpctl.SyclQueue("opencl") except dpctl.SyclQueueCreationError: pytest.skip("No OpenCL queue is available") - kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) _check_multi_kernel_program(kb) @@ -217,7 +217,7 @@ def test_create_kernel_bundle_from_spirv_ocl(): spirv_file = get_spirv_abspath("multi_kernel.spv") with open(spirv_file, "rb") as fin: spirv = fin.read() - kb = dpctl_prog.create_kernel_bundle_from_spirv(q, spirv) + kb = dpc.create_kernel_bundle_from_spirv(q, spirv) _check_multi_kernel_program(kb) @@ -229,7 +229,7 @@ def test_create_kernel_bundle_from_spirv_l0(): spirv_file = get_spirv_abspath("multi_kernel.spv") with open(spirv_file, "rb") as fin: spirv = fin.read() - kb = dpctl_prog.create_kernel_bundle_from_spirv(q, spirv) + kb = dpc.create_kernel_bundle_from_spirv(q, spirv) _check_multi_kernel_program(kb) @@ -250,7 +250,7 @@ def test_create_kernel_bundle_from_source_l0(): size_t index = get_global_id(0); \ c[index] = a[index] + d*b[index]; \ }" - kb = dpctl_prog.create_kernel_bundle_from_source(q, oclSrc) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) _check_multi_kernel_program(kb) @@ -262,8 +262,8 @@ def test_create_kernel_bundle_from_invalid_src_ocl(): invalid_oclSrc = " \ kernel void add( \ }" - with pytest.raises(dpctl_prog.SyclKernelBundleCompilationError): - dpctl_prog.create_kernel_bundle_from_source(q, invalid_oclSrc) + with pytest.raises(dpc.SyclKernelBundleCompilationError): + dpc.create_kernel_bundle_from_source(q, invalid_oclSrc) def test_create_kernel_bundle_with_spec_const(): @@ -273,15 +273,13 @@ def test_create_kernel_bundle_with_spec_const(): pytest.skip("Could not create default queue") spec_id = 0 - sp = dpctl_prog.SpecializationConstant(spec_id, "i4", 42) + sp = dpc.SpecializationConstant(spec_id, "i4", 42) spirv_file = get_spirv_abspath("specialization_constant_kernel.spv") with open(spirv_file, "br") as spv: spv_bytes = spv.read() - kb = dpctl_prog.create_kernel_bundle_from_spirv( - q, spv_bytes, specializations=[sp] - ) + kb = dpc.create_kernel_bundle_from_spirv(q, spv_bytes, specializations=[sp]) kernel = kb.get_sycl_kernel("_ZTS20BasicSpecConstKernel") n = 128 @@ -311,15 +309,15 @@ def test_create_kernel_bundle_with_composite_spec_const(): # composite specialization constants are separated into individual # specialization constants with unique spec_ids - sp1 = dpctl_prog.SpecializationConstant(0, "i4", 10) - sp2 = dpctl_prog.SpecializationConstant(1, "f4", 2.5) - sp3 = dpctl_prog.SpecializationConstant(2, "?", 1) + sp1 = dpc.SpecializationConstant(0, "i4", 10) + sp2 = dpc.SpecializationConstant(1, "f4", 2.5) + sp3 = dpc.SpecializationConstant(2, "?", 1) spirv_file = get_spirv_abspath("specialization_constant_composite.spv") with open(spirv_file, "br") as spv: spv_bytes = spv.read() - kb = dpctl_prog.create_kernel_bundle_from_spirv( + kb = dpc.create_kernel_bundle_from_spirv( q, spv_bytes, specializations=[sp1, sp2, sp3] ) kernel = kb.get_sycl_kernel("_ZTS21StructSpecConstKernel") @@ -345,12 +343,12 @@ def test_create_kernel_bundle_with_composite_spec_const(): def test_specialization_constant_addressof_ref(): - sp = dpctl_prog.SpecializationConstant(0, "i4", 42) + sp = dpc.SpecializationConstant(0, "i4", 42) ref = sp.addressof_ref() assert type(ref) is int assert ref != 0 - sp2 = dpctl_prog.SpecializationConstant(1, "f4", 3.14) + sp2 = dpc.SpecializationConstant(1, "f4", 3.14) ref2 = sp2.addressof_ref() assert type(ref2) is int assert ref2 != 0 @@ -404,10 +402,10 @@ def test_spirv_specializations_parser_invalid_spirv(): def test_spec_const_equality(): - sp1 = dpctl_prog.SpecializationConstant(0, "i4", 42) - sp2 = dpctl_prog.SpecializationConstant(0, "i4", 42) - sp3 = dpctl_prog.SpecializationConstant(1, "i4", 42) - sp4 = dpctl_prog.SpecializationConstant(0, "f4", 42.0) + sp1 = dpc.SpecializationConstant(0, "i4", 42) + sp2 = dpc.SpecializationConstant(0, "i4", 42) + sp3 = dpc.SpecializationConstant(1, "i4", 42) + sp4 = dpc.SpecializationConstant(0, "f4", 42.0) assert sp1 == sp2 assert sp1 != sp3 diff --git a/dpctl/tests/test_work_group_memory.py b/dpctl/tests/test_work_group_memory.py index 942fd15e49..1c648dd2e2 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() - kb = dpctl.program.create_kernel_bundle_from_spirv(q, spv_bytes) + kb = dpctl.compiler.create_kernel_bundle_from_spirv(q, spv_bytes) kernel = kb.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 46eb9a906e..f542ab8c03 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") - kb = dpctl.program.create_kernel_bundle_from_source(q, ocl_kernel_src) + kb = dpctl.compiler.create_kernel_bundle_from_source(q, ocl_kernel_src) kernel = kb.get_sycl_kernel("local_mem_kernel") local_size = 16 global_size = local_size * 8 diff --git a/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt index 5d8129581e..5bb2a98c6f 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt +++ b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt @@ -1,7 +1,7 @@ cmake_minimum_required(VERSION 3.21...3.27 FATAL_ERROR) project(use_queue_device VERSION 0.1 LANGUAGES CXX - DESCRIPTION "Example of using dpctl.program.SyclKernel <-> sycl::kernel type casting") + DESCRIPTION "Example of using dpctl.compiler.SyclKernel <-> sycl::kernel type casting") set(DPCTL_CMAKE_MODULES_PATH "${CMAKE_SOURCE_DIR}/../../../cmake") set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${DPCTL_CMAKE_MODULES_PATH}) diff --git a/examples/pybind11/use_dpctl_sycl_kernel/README.md b/examples/pybind11/use_dpctl_sycl_kernel/README.md index 77aa57bf6e..34fa8d96ad 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/README.md +++ b/examples/pybind11/use_dpctl_sycl_kernel/README.md @@ -3,7 +3,7 @@ ## Description This extension demonstrates how you can use dpctl Python types, -such as ``dpctl.SyclQueue`` and ``dpctl.program.SyclKernel``, in +such as ``dpctl.SyclQueue`` and ``dpctl.compiler.SyclKernel``, in Pybind11 extensions. diff --git a/examples/pybind11/use_dpctl_sycl_kernel/example.py b/examples/pybind11/use_dpctl_sycl_kernel/example.py index 272f32cdbf..10d6b39e12 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/example.py +++ b/examples/pybind11/use_dpctl_sycl_kernel/example.py @@ -20,8 +20,8 @@ import use_kernel as eg import dpctl -import dpctl.memory as dpmem -import dpctl.program as dppr +import dpctl.compiler as dpc +import dpctl.memory as dpm # create execution queue, targeting default selected device q = dpctl.SyclQueue() @@ -31,7 +31,7 @@ il = fh.read() # Build the program for the selected device -pr = dppr.create_kernel_bundle_from_spirv(q, il, "") +pr = dpc.create_kernel_bundle_from_spirv(q, il, "") assert pr.has_sycl_kernel("double_it") # Retrieve the kernel from the problem @@ -41,8 +41,8 @@ # Construct the argument, and allocate memory for the result x = np.arange(0, stop=13, step=1, dtype="i4") y = np.empty_like(x) -x_dev = dpmem.MemoryUSMDevice(x.nbytes, queue=q) -y_dev = dpmem.MemoryUSMDevice(y.nbytes, queue=q) +x_dev = dpm.MemoryUSMDevice(x.nbytes, queue=q) +y_dev = dpm.MemoryUSMDevice(y.nbytes, queue=q) # Copy input data to the device q.memcpy(dest=x_dev, src=x, count=x.nbytes) 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 f061697442..aab82e23ed 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 @@ -23,8 +23,8 @@ import use_kernel as uk import dpctl -import dpctl.memory as dpmem -import dpctl.program as dppr +import dpctl.compiler as dpc +import dpctl.memory as dpm def _get_spv_path(): @@ -45,7 +45,7 @@ def test_kernel_can_be_found(): q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: pytest.skip("Could not create default queue") - kb = dppr.create_kernel_bundle_from_spirv(q, il, "") + kb = dpc.create_kernel_bundle_from_spirv(q, il, "") assert kb.has_sycl_kernel("double_it") @@ -57,15 +57,15 @@ def test_kernel_submit_through_extension(): q = dpctl.SyclQueue() except dpctl.SyclQueueCreationError: pytest.skip("Could not create default queue") - kb = dppr.create_kernel_bundle_from_spirv(q, il, "") + kb = dpc.create_kernel_bundle_from_spirv(q, il, "") krn = kb.get_sycl_kernel("double_it") assert krn.num_args == 2 x = np.arange(0, stop=13, step=1, dtype="i4") y = np.empty_like(x) - x_usm = dpmem.MemoryUSMDevice(x.nbytes, queue=q) - y_usm = dpmem.MemoryUSMDevice(y.nbytes, queue=q) + x_usm = dpm.MemoryUSMDevice(x.nbytes, queue=q) + y_usm = dpm.MemoryUSMDevice(y.nbytes, queue=q) ev = q.memcpy_async(dest=x_usm, src=x, count=x_usm.nbytes) diff --git a/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py index 28da2eaa4c..d2dce3de12 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py +++ b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py @@ -27,7 +27,7 @@ SYCL entities. dpctl provides type casters that bind ``sycl::kernel`` to -`dpctl.program.SyclKernel`, ``sycl::device`` to `dpctl.SyclDevice`, etc. +`dpctl.compiler.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/setup.py b/setup.py index 17cb7a9492..d4c9987b1d 100644 --- a/setup.py +++ b/setup.py @@ -25,9 +25,10 @@ url="https://github.com/IntelPython/dpctl", packages=[ "dpctl", + "dpctl.compiler", + "dpctl.compiler.utils", "dpctl.memory", "dpctl.program", - "dpctl.program.utils", "dpctl.utils", ], package_data={ @@ -44,10 +45,10 @@ "include/syclinterface/Support/*.h", "include/dpctl/_sycl*.h", "include/dpctl/memory/_memory*.h", - "include/dpctl/program/_program*.h", + "include/dpctl/compiler/_compiler*.h", "*.pxd", + "compiler/*.pxd", "memory/*.pxd", - "program/*.pxd", ] }, include_package_data=False,