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 c033895c19..9ac87f4eb1 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -10,6 +10,9 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * `dpctl.SyclQueue.copy` and `dpctl.SyclQueue.copy_async` methods [gh-2273](https://github.com/IntelPython/dpctl/pull/2273) * Added a number of `sycl::device` info queries to `dpctl.SyclDevice` [gh-2324](https://github.com/IntelPython/dpctl/pull/2324) +### 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) @@ -17,6 +20,10 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### 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) +* Deprecated `DPCTL_ENABLE_L0_PROGRAM_CREATION` CMake option in favor of `DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION` [gh-2317](https://github.com/IntelPython/dpctl/pull/2317) + ### Maintenance * Updated pybind11 version used by `dpctl` and examples [gh-2357](https://github.com/IntelPython/dpctl/pull/2357) 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/docs/doc_sources/contributor_guides/building.rst b/docs/doc_sources/contributor_guides/building.rst index adb0dbd83a..3975941f56 100644 --- a/docs/doc_sources/contributor_guides/building.rst +++ b/docs/doc_sources/contributor_guides/building.rst @@ -250,7 +250,7 @@ library. -DCMAKE_CXX_COMPILER=icpx \ -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ - -DDPCTL_ENABLE_L0_PROGRAM_CREATION=ON \ + -DDPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION=ON \ -DDPCTL_BUILD_CAPI_TESTS=ON \ -DDPCTL_GENERATE_COVERAGE=ON \ .. 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 42d61207de..11c4a3281f 100644 --- a/dpctl/__init__.py +++ b/dpctl/__init__.py @@ -135,6 +135,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..ecfc85596f 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -88,7 +88,7 @@ class dpctl_capi DPCTLSyclQueueRef, PyObject *); - // program + // compiler DPCTLSyclKernelRef (*SyclKernel_GetKernelRef_)(PySyclKernelObject *); PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef, const char *); @@ -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_program.py b/dpctl/tests/test_sycl_compiler.py similarity index 83% rename from dpctl/tests/test_sycl_program.py rename to dpctl/tests/test_sycl_compiler.py index bbf16df98c..d7364d8de1 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_compiler.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): @@ -32,14 +32,14 @@ def get_spirv_abspath(fn): return spirv_file -def _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(sycl_prog): +def _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(kb): """Checks Cython-generated C-API function - `SyclKernelBundle_GetKernelBundleRef` defined in _program.pyx""" + `SyclKernelBundle_GetKernelBundleRef` defined in _compiler.pyx""" import ctypes import sys - assert type(sycl_prog) is dpctl_prog.SyclKernelBundle - mod = sys.modules[sycl_prog.__class__.__module__] + assert type(kb) is dpc.SyclKernelBundle + mod = sys.modules[kb.__class__.__module__] # get capsule storing SyclKernelBundle_GetKernelBundleRef function ptr kb_ref_fn_cap = mod.__pyx_capi__["SyclKernelBundle_GetKernelBundleRef"] # construct Python callable to invoke "SyclKernelBundle_GetKernelBundleRef" @@ -54,45 +54,45 @@ def _check_cpython_api_SyclKernelBundle_GetKernelBundleRef(sycl_prog): 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) + r2 = kb.addressof_ref() + r1 = get_kernel_bundle_ref_fn(kb) assert r1 == r2 -def _check_cpython_api_SyclKernelBundle_Make(sycl_prog): +def _check_cpython_api_SyclKernelBundle_Make(kb): """Checks Cython-generated C-API function - `SyclKernelBundle_Make` defined in _program.pyx""" + `SyclKernelBundle_Make` defined in _compiler.pyx""" import ctypes import sys - assert type(sycl_prog) is dpctl_prog.SyclKernelBundle - mod = sys.modules[sycl_prog.__class__.__module__] + assert type(kb) is dpc.SyclKernelBundle + mod = sys.modules[kb.__class__.__module__] # get capsule storing SyclKernelBundle_Make function ptr - make_prog_fn_cap = mod.__pyx_capi__["SyclKernelBundle_Make"] + make_kb_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, + make_kb_fn_ptr = cap_ptr_fn( + make_kb_fn_cap, b"struct PySyclKernelBundleObject *(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) + make_kb_fn = callable_maker(make_kb_fn_ptr) - p2 = make_prog_fn(sycl_prog.addressof_ref()) - assert p2.has_sycl_kernel("add") - assert p2.has_sycl_kernel("axpy") + kb2 = make_kb_fn(kb.addressof_ref()) + assert kb2.has_sycl_kernel("add") + assert kb2.has_sycl_kernel("axpy") def _check_cpython_api_SyclKernel_GetKernelRef(krn): """Checks Cython-generated C-API function - `SyclKernel_GetKernelRef` defined in _program.pyx""" + `SyclKernel_GetKernelRef` defined in _compiler.pyx""" 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"] @@ -114,11 +114,11 @@ def _check_cpython_api_SyclKernel_GetKernelRef(krn): def _check_cpython_api_SyclKernel_Make(krn): """Checks Cython-generated C-API function - `SyclKernel_Make` defined in _program.pyx""" + `SyclKernel_Make` defined in _compiler.pyx""" 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"] @@ -149,8 +149,8 @@ def _check_cpython_api_SyclKernel_Make(krn): assert krn.work_group_size == k3.work_group_size -def _check_multi_kernel_program(kb): - assert type(kb) is dpctl_prog.SyclKernelBundle +def _check_multi_kernel_bundle(kb): + assert type(kb) is dpc.SyclKernelBundle assert type(kb.addressof_ref()) is int assert kb.has_sycl_kernel("add") @@ -205,8 +205,8 @@ 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) - _check_multi_kernel_program(kb) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) + _check_multi_kernel_bundle(kb) def test_create_kernel_bundle_from_spirv_ocl(): @@ -217,8 +217,8 @@ 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) - _check_multi_kernel_program(kb) + kb = dpc.create_kernel_bundle_from_spirv(q, spirv) + _check_multi_kernel_bundle(kb) def test_create_kernel_bundle_from_spirv_l0(): @@ -229,8 +229,8 @@ 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) - _check_multi_kernel_program(kb) + kb = dpc.create_kernel_bundle_from_spirv(q, spirv) + _check_multi_kernel_bundle(kb) @pytest.mark.xfail( @@ -250,8 +250,8 @@ 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) - _check_multi_kernel_program(kb) + kb = dpc.create_kernel_bundle_from_source(q, oclSrc) + _check_multi_kernel_bundle(kb) def test_create_kernel_bundle_from_invalid_src_ocl(): @@ -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_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_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 6d4ec45893..f48b8c83cf 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..e0f7276954 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() @@ -30,19 +30,19 @@ with open("resource/double_it.spv", "br") as fh: il = fh.read() -# Build the program for the selected device -pr = dppr.create_kernel_bundle_from_spirv(q, il, "") -assert pr.has_sycl_kernel("double_it") +# Build the kernel bundle for the selected device +kb = dpc.create_kernel_bundle_from_spirv(q, il, "") +assert kb.has_sycl_kernel("double_it") -# Retrieve the kernel from the problem -krn = pr.get_sycl_kernel("double_it") +# Retrieve the kernel from the kernel bundle +krn = kb.get_sycl_kernel("double_it") assert krn.num_args == 2 # 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/libsyclinterface/CMakeLists.txt b/libsyclinterface/CMakeLists.txt index e71080602a..3e7e73cdd4 100644 --- a/libsyclinterface/CMakeLists.txt +++ b/libsyclinterface/CMakeLists.txt @@ -15,12 +15,26 @@ if(NOT DEFINED IntelSYCL_FOUND OR NOT IntelSYCL_FOUND) find_package(IntelSYCL REQUIRED) endif() -# Option to turn on support for creating Level Zero interoperability programs -# from a SPIR-V binary file. -option(DPCTL_ENABLE_L0_PROGRAM_CREATION - "Enable Level Zero Program creation from SPIR-V" +# Option to turn on support for creating Level Zero interoperability kernel +# bundles from a SPIR-V binary file. +option(DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION + "Enable Level Zero kernel bundle creation from SPIR-V" ON ) +if(DEFINED DPCTL_ENABLE_L0_PROGRAM_CREATION) + message(DEPRECATION + "DPCTL_ENABLE_L0_PROGRAM_CREATION is deprecated and will be removed in " + "a future release. Use DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION instead." + ) + set(DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION + ${DPCTL_ENABLE_L0_PROGRAM_CREATION} + CACHE BOOL + "Enable Level Zero kernel bundle creation from SPIR-V" + FORCE + ) + unset(DPCTL_ENABLE_L0_PROGRAM_CREATION) + unset(DPCTL_ENABLE_L0_PROGRAM_CREATION CACHE) +endif() # Option to generate code coverage report using llvm-cov and lcov. option(DPCTL_GENERATE_COVERAGE "Build dpctl C API with coverage instrumentation" @@ -59,8 +73,8 @@ else() find_package(IntelSyclCompiler REQUIRED) endif() -if(DPCTL_ENABLE_L0_PROGRAM_CREATION) - set(DPCTL_ENABLE_L0_PROGRAM_CREATION 1) +if(DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION) + set(DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION 1) if(DPCTL_LEVEL_ZERO_INCLUDE_DIR) set(LEVEL_ZERO_INCLUDE_DIR ${DPCTL_LEVEL_ZERO_INCLUDE_DIR}) else() @@ -371,7 +385,7 @@ if (build_so_version) ) endif() -if(DPCTL_ENABLE_L0_PROGRAM_CREATION) +if(DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION) target_include_directories(DPCTLSyclInterface SYSTEM PRIVATE ${LEVEL_ZERO_INCLUDE_DIR} diff --git a/libsyclinterface/dbg_build.sh b/libsyclinterface/dbg_build.sh index c090bc700c..2da0d29fb5 100755 --- a/libsyclinterface/dbg_build.sh +++ b/libsyclinterface/dbg_build.sh @@ -20,7 +20,7 @@ cmake \ -DCMAKE_CXX_FLAGS=-fsycl \ -DCMAKE_INSTALL_PREFIX="${INSTALL_PREFIX}" \ -DCMAKE_PREFIX_PATH="${INSTALL_PREFIX}" \ - -DDPCTL_ENABLE_L0_PROGRAM_CREATION=ON \ + -DDPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION=ON \ -DDPCTL_BUILD_CAPI_TESTS=ON \ -DDPCTL_GENERATE_COVERAGE=OFF \ .. diff --git a/libsyclinterface/dbg_build_custom.sh b/libsyclinterface/dbg_build_custom.sh index 4c30cca1ce..61dc103bf8 100755 --- a/libsyclinterface/dbg_build_custom.sh +++ b/libsyclinterface/dbg_build_custom.sh @@ -17,7 +17,7 @@ cmake \ -DCMAKE_PREFIX_PATH="${INSTALL_PREFIX}" \ -DDPCTL_CUSTOM_DPCPP_INSTALL_DIR="${DPCPP_HOME}" \ -DCMAKE_LINKER:PATH="${DPCPP_HOME}/bin/lld" \ - -DDPCTL_ENABLE_L0_PROGRAM_CREATION="${USE_LO_HEADERS}" \ + -DDPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION="${USE_LO_HEADERS}" \ -DDPCTL_BUILD_CAPI_TESTS=ON \ -DDPCTL_GENERATE_COVERAGE=ON \ .. diff --git a/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in b/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in index 97e28d1219..9900742234 100644 --- a/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in +++ b/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in @@ -25,8 +25,9 @@ #pragma once -/* Defined when dpctl was built with level zero program creation enabled. */ -#cmakedefine DPCTL_ENABLE_L0_PROGRAM_CREATION 1 +/* Defined when dpctl was built with Level Zero kernel bundle creation + * enabled. */ +#cmakedefine DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION 1 #define __SYCL_COMPILER_VERSION_REQUIRED 20221201L diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index ec94e5a077..1733ecef5b 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -38,7 +38,7 @@ #include /* Sycl headers */ #include -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +#ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION // Note: include ze_api.h before level_zero.hpp. Make sure clang-format does // not reorder the includes. // clang-format off @@ -63,7 +63,7 @@ static const int clLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; static const char *clLoaderName = "OpenCL.dll"; static const int clLibLoadFlags = 0; #else -#error "OpenCL program compilation is unavailable for this platform" +#error "OpenCL kernel bundle compilation is unavailable for this platform" #endif #define CodeStringSuffix(code) \ @@ -408,7 +408,7 @@ _GetKernel_ocl_impl(const kernel_bundle &kb, } } -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +#ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION #ifdef __linux__ static const char *zeLoaderName = DPCTL_LIBZE_LOADER_FILENAME; @@ -417,7 +417,7 @@ static const int zeLibLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; static const char *zeLoaderName = "ze_loader.dll"; static const int zeLibLoadFlags = 0; #else -#error "Level Zero program compilation is unavailable for this platform" +#error "Level Zero kernel bundle compilation is unavailable for this platform" #endif static constexpr sycl::backend ze_be = sycl::backend::ext_oneapi_level_zero; @@ -676,7 +676,7 @@ bool _HasKernel_ze_impl(const kernel_bundle &kb, return false; } -#endif /* #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION */ +#endif /* #ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION */ } /* end of anonymous namespace */ @@ -691,19 +691,19 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, { DPCTLSyclKernelBundleRef KBRef = nullptr; if (!CtxRef) { - error_handler("Cannot create program from SPIR-V as the supplied SYCL " - "context is NULL.", + error_handler("Cannot create kernel bundle from SPIR-V as the supplied " + "SYCL context is NULL.", __FILE__, __func__, __LINE__); return KBRef; } if (!DevRef) { - error_handler("Cannot create program from SPIR-V as the supplied SYCL " - "device is NULL.", + error_handler("Cannot create kernel bundle from SPIR-V as the supplied " + "SYCL device is NULL.", __FILE__, __func__, __LINE__); return KBRef; } if ((!IL) || (length == 0)) { - error_handler("Cannot create program from null SPIR-V buffer.", + error_handler("Cannot create kernel bundle from null SPIR-V buffer.", __FILE__, __func__, __LINE__); return KBRef; } @@ -720,7 +720,7 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, SpecConsts); break; case backend::ext_oneapi_level_zero: -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +#ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION KBRef = _CreateKernelBundleWithIL_ze_impl( *SyclCtx, *SyclDev, IL, length, CompileOpts, NumSpecConsts, SpecConsts); @@ -807,7 +807,7 @@ DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, case sycl::backend::opencl: return _GetKernel_ocl_impl(*SyclKB, KernelName); case sycl::backend::ext_oneapi_level_zero: -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +#ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION return _GetKernel_ze_impl(*SyclKB, KernelName); #endif default: @@ -837,7 +837,7 @@ bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, case sycl::backend::opencl: return _HasKernel_ocl_impl(*SyclKB, KernelName); case sycl::backend::ext_oneapi_level_zero: -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +#ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION return _HasKernel_ze_impl(*SyclKB, KernelName); #endif default: diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index 7917157c28..aca6a1b307 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -283,7 +283,7 @@ INSTANTIATE_TEST_SUITE_P(KernelBundleCreationFromSpirv, "opencl:gpu", "opencl:cpu", "opencl:gpu:0", -#ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION +#ifdef DPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION "level_zero", "level_zero:gpu", #endif diff --git a/scripts/_build_helper.py b/scripts/_build_helper.py index d41567598b..dfbc86f46d 100644 --- a/scripts/_build_helper.py +++ b/scripts/_build_helper.py @@ -97,7 +97,8 @@ def make_cmake_args( args = [ f"-DCMAKE_C_COMPILER:PATH={c_compiler}" if c_compiler else "", f"-DCMAKE_CXX_COMPILER:PATH={cxx_compiler}" if cxx_compiler else "", - f"-DDPCTL_ENABLE_L0_PROGRAM_CREATION={'ON' if level_zero else 'OFF'}", + "-DDPCTL_ENABLE_L0_KERNEL_BUNDLE_CREATION=" + f"{'ON' if level_zero else 'OFF'}", f"-DDPCTL_ENABLE_GLOG:BOOL={'ON' if glog else 'OFF'}", ] 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,