From 963d2fc9bb9b62f32458cd19c78c35624ad16101 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:16:38 -0700 Subject: [PATCH 01/12] deprecate dpctl program in favor of dpctl compiler --- .flake8 | 2 +- dpctl/CMakeLists.txt | 1 + dpctl/__init__.py | 1 + dpctl/apis/include/dpctl4pybind11.hpp | 14 ++-- dpctl/{program => compiler}/CMakeLists.txt | 2 +- dpctl/compiler/__init__.pxd | 25 ++++++ dpctl/compiler/__init__.py | 76 +++++++++++++++++++ dpctl/{program => compiler}/_program.pxd | 0 dpctl/{program => compiler}/_program.pyx | 0 dpctl/compiler/utils/__init__.py | 25 ++++++ dpctl/{program => compiler}/utils/_utils.py | 2 +- dpctl/program/__init__.pxd | 8 +- dpctl/program/__init__.py | 31 +++++--- dpctl/program/utils/__init__.py | 14 +++- dpctl/tests/test_sycl_context.py | 10 +-- dpctl/tests/test_sycl_event.py | 16 ++-- dpctl/tests/test_sycl_kernel_submit.py | 14 ++-- dpctl/tests/test_sycl_program.py | 40 +++++----- .../pybind11/use_dpctl_sycl_kernel/README.md | 2 +- .../pybind11/use_dpctl_sycl_kernel/example.py | 10 +-- .../tests/test_user_kernel.py | 12 +-- .../use_kernel/__init__.py | 2 +- setup.py | 4 + 23 files changed, 230 insertions(+), 81 deletions(-) rename dpctl/{program => compiler}/CMakeLists.txt (72%) create mode 100644 dpctl/compiler/__init__.pxd create mode 100644 dpctl/compiler/__init__.py rename dpctl/{program => compiler}/_program.pxd (100%) rename dpctl/{program => compiler}/_program.pyx (100%) create mode 100644 dpctl/compiler/utils/__init__.py rename dpctl/{program => compiler}/utils/_utils.py (99%) diff --git a/.flake8 b/.flake8 index 1453996659..278bc6ae29 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/_program.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/dpctl/CMakeLists.txt b/dpctl/CMakeLists.txt index a24c7443f9..3b259927b9 100644 --- a/dpctl/CMakeLists.txt +++ b/dpctl/CMakeLists.txt @@ -203,6 +203,7 @@ 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(compiler) add_subdirectory(program) 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/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/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..e6feca320c 100644 --- a/dpctl/program/CMakeLists.txt +++ b/dpctl/compiler/CMakeLists.txt @@ -2,6 +2,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/compiler/__init__.pxd b/dpctl/compiler/__init__.pxd new file mode 100644 index 0000000000..9b976be265 --- /dev/null +++ b/dpctl/compiler/__init__.pxd @@ -0,0 +1,25 @@ +# 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. + +# distutils: language = c++ +# cython: language_level=3 + +"""Declares the extension types and functions for the Cython API +implemented in dpctl.compiler._program.pyx. +""" + + +from dpctl.compiler._program cimport * diff --git a/dpctl/compiler/__init__.py b/dpctl/compiler/__init__.py new file mode 100644 index 0000000000..20cbc24453 --- /dev/null +++ b/dpctl/compiler/__init__.py @@ -0,0 +1,76 @@ +# 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._program import ( + SpecializationConstant, + SyclKernel, + SyclKernelBundle, + SyclKernelBundleCompilationError, + create_kernel_bundle_from_source, + create_kernel_bundle_from_spirv, + create_program_from_source, + create_program_from_spirv, +) + +__all__ = [ + "create_kernel_bundle_from_source", + "create_kernel_bundle_from_spirv", + "create_program_from_source", + "create_program_from_spirv", + "SyclKernel", + "SyclKernelBundle", + "SyclKernelBundleCompilationError", + "SyclProgram", + "SyclProgramCompilationError", + "SpecializationConstant", +] + +# add submodules +__all__ += [ + "utils", +] + + +def __getattr__(name): + if name == "SyclProgram": + from warnings import warn + + warn( + "dpctl.compiler.SyclProgram is deprecated and will be removed in a " + "future release. Use dpctl.compiler.SyclKernelBundle instead.", + DeprecationWarning, + stacklevel=2, + ) + return SyclKernelBundle + if name == "SyclProgramCompilationError": + from warnings import warn + + warn( + "dpctl.compiler.SyclProgramCompilationError is deprecated and will " + "be removed in a future release. Use " + "dpctl.compiler.SyclKernelBundleCompilationError instead.", + DeprecationWarning, + stacklevel=2, + ) + return SyclKernelBundleCompilationError + raise AttributeError(f"module {__name__} has no attribute {name}") diff --git a/dpctl/program/_program.pxd b/dpctl/compiler/_program.pxd similarity index 100% rename from dpctl/program/_program.pxd rename to dpctl/compiler/_program.pxd diff --git a/dpctl/program/_program.pyx b/dpctl/compiler/_program.pyx similarity index 100% rename from dpctl/program/_program.pyx rename to dpctl/compiler/_program.pyx diff --git a/dpctl/compiler/utils/__init__.py b/dpctl/compiler/utils/__init__.py new file mode 100644 index 0000000000..1359230360 --- /dev/null +++ b/dpctl/compiler/utils/__init__.py @@ -0,0 +1,25 @@ +# 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. + +""" +A collection of utility functions for dpctl.compiler module. +""" + +from dpctl.compiler.utils._utils import parse_spirv_specializations + +__all__ = [ + "parse_spirv_specializations", +] 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__.pxd b/dpctl/program/__init__.pxd index 38b71beebf..688b002467 100644 --- a/dpctl/program/__init__.pxd +++ b/dpctl/program/__init__.pxd @@ -17,9 +17,11 @@ # distutils: language = c++ # cython: language_level=3 -"""Declares the extension types and functions for the Cython API -implemented in dpctl.program._program.pyx. +""" +Declares the extension types and functions for the Cython API +implemented in dpctl.compiler._program.pyx (deprecated, use dpctl.compiler +instead). """ -from dpctl.program._program cimport * +from dpctl.compiler._program cimport * diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index 4eedbee7a8..e1ef0c5809 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -15,14 +15,18 @@ # 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 ( +import warnings + +from dpctl.compiler import ( SpecializationConstant, SyclKernel, SyclKernelBundle, @@ -51,25 +55,28 @@ "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/utils/__init__.py b/dpctl/program/utils/__init__.py index 34c480838d..fdf5f8e67f 100644 --- a/dpctl/program/utils/__init__.py +++ b/dpctl/program/utils/__init__.py @@ -15,12 +15,22 @@ # limitations under the License. """ -A collection of utility functions for dpctl.program module. +A collection of utility functions for dpctl.program module (deprecated, use +dpctl.compiler.utils instead). """ -from ._utils import SpecializationConstantInfo, parse_spirv_specializations +import warnings + +from dpctl.compiler.utils import SpecializationConstantInfo, parse_spirv_specializations __all__ = [ "parse_spirv_specializations", "SpecializationConstantInfo", ] + +warnings.warn( + "dpctl.program.utils is deprecated and will be removed in a future release." + " Use dpctl.compiler.utils instead.", + DeprecationWarning, + stacklevel=2, +) 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..4cc8d8a4fc 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") 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..169ccfc8f0 100644 --- a/setup.py +++ b/setup.py @@ -25,6 +25,8 @@ url="https://github.com/IntelPython/dpctl", packages=[ "dpctl", + "dpctl.compiler", + "dpctl.compiler.utils", "dpctl.memory", "dpctl.program", "dpctl.program.utils", @@ -44,8 +46,10 @@ "include/syclinterface/Support/*.h", "include/dpctl/_sycl*.h", "include/dpctl/memory/_memory*.h", + "include/dpctl/compiler/_program*.h", "include/dpctl/program/_program*.h", "*.pxd", + "compiler/*.pxd", "memory/*.pxd", "program/*.pxd", ] From 088917a0d022e65135ea645bbd0a9fd5fe3d6783 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:18:03 -0700 Subject: [PATCH 02/12] update docs to contain dpctl.compiler --- .../api_reference/dpctl/compiler.rst | 39 +++++++++++++++++++ .../doc_sources/api_reference/dpctl/index.rst | 4 +- .../api_reference/dpctl/program.rst | 7 +++- docs/doc_sources/api_reference/index.rst | 4 +- 4 files changed, 50 insertions(+), 4 deletions(-) create mode 100644 docs/doc_sources/api_reference/dpctl/compiler.rst 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..965638aa78 --- /dev/null +++ b/docs/doc_sources/api_reference/dpctl/compiler.rst @@ -0,0 +1,39 @@ +.. _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 + create_program_from_source + create_program_from_spirv + +.. autosummary:: + :toctree: generated + :nosignatures: + + SyclKernelBundle + SyclKernel + +.. 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..8e86a236bb 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -1,7 +1,10 @@ .. _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 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 From 5c09bfcc262a8e00022c21f80f9cfb60af80e3a0 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:19:51 -0700 Subject: [PATCH 03/12] remove star import in dpctl.compiler --- dpctl/compiler/__init__.pxd | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/dpctl/compiler/__init__.pxd b/dpctl/compiler/__init__.pxd index 9b976be265..f9da512601 100644 --- a/dpctl/compiler/__init__.pxd +++ b/dpctl/compiler/__init__.pxd @@ -22,4 +22,11 @@ implemented in dpctl.compiler._program.pyx. """ -from dpctl.compiler._program cimport * +from dpctl.compiler._program cimport ( + SyclKernel, + SyclKernelBundle, + create_kernel_bundle_from_source, + create_kernel_bundle_from_spirv, + create_program_from_source, + create_program_from_spirv, +) From 061cbb0a55dd2b83109438b250a9a1cad5c030a4 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:32:26 -0700 Subject: [PATCH 04/12] move deprecated APIs into program submodule --- .../api_reference/dpctl/compiler.rst | 3 +- .../api_reference/dpctl/program.rst | 29 +-------- dpctl/compiler/__init__.py | 31 ---------- dpctl/compiler/_program.pyx | 32 ---------- dpctl/program/__init__.py | 14 +---- dpctl/program/_program.py | 60 +++++++++++++++++++ dpctl/program/utils/__init__.py | 36 ----------- setup.py | 1 - 8 files changed, 65 insertions(+), 141 deletions(-) create mode 100644 dpctl/program/_program.py delete mode 100644 dpctl/program/utils/__init__.py diff --git a/docs/doc_sources/api_reference/dpctl/compiler.rst b/docs/doc_sources/api_reference/dpctl/compiler.rst index 965638aa78..c887ecf562 100644 --- a/docs/doc_sources/api_reference/dpctl/compiler.rst +++ b/docs/doc_sources/api_reference/dpctl/compiler.rst @@ -22,8 +22,6 @@ execution via :py:meth:`dpctl.SyclQueue.submit`. create_kernel_bundle_from_source create_kernel_bundle_from_spirv - create_program_from_source - create_program_from_spirv .. autosummary:: :toctree: generated @@ -31,6 +29,7 @@ execution via :py:meth:`dpctl.SyclQueue.submit`. SyclKernelBundle SyclKernel + SpecializationConstant .. autosummary:: :toctree: generated diff --git a/docs/doc_sources/api_reference/dpctl/program.rst b/docs/doc_sources/api_reference/dpctl/program.rst index 8e86a236bb..619947fd6c 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -23,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/dpctl/compiler/__init__.py b/dpctl/compiler/__init__.py index 20cbc24453..d6d3f8f0d2 100644 --- a/dpctl/compiler/__init__.py +++ b/dpctl/compiler/__init__.py @@ -28,20 +28,14 @@ SyclKernelBundleCompilationError, create_kernel_bundle_from_source, create_kernel_bundle_from_spirv, - create_program_from_source, - create_program_from_spirv, ) __all__ = [ "create_kernel_bundle_from_source", "create_kernel_bundle_from_spirv", - "create_program_from_source", - "create_program_from_spirv", "SyclKernel", "SyclKernelBundle", "SyclKernelBundleCompilationError", - "SyclProgram", - "SyclProgramCompilationError", "SpecializationConstant", ] @@ -49,28 +43,3 @@ __all__ += [ "utils", ] - - -def __getattr__(name): - if name == "SyclProgram": - from warnings import warn - - warn( - "dpctl.compiler.SyclProgram is deprecated and will be removed in a " - "future release. Use dpctl.compiler.SyclKernelBundle instead.", - DeprecationWarning, - stacklevel=2, - ) - return SyclKernelBundle - if name == "SyclProgramCompilationError": - from warnings import warn - - warn( - "dpctl.compiler.SyclProgramCompilationError is deprecated and will " - "be removed in a future release. Use " - "dpctl.compiler.SyclKernelBundleCompilationError instead.", - DeprecationWarning, - stacklevel=2, - ) - return SyclKernelBundleCompilationError - raise AttributeError(f"module {__name__} has no attribute {name}") diff --git a/dpctl/compiler/_program.pyx b/dpctl/compiler/_program.pyx index 1c9227ec02..1f3d154cca 100644 --- a/dpctl/compiler/_program.pyx +++ b/dpctl/compiler/_program.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, @@ -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/__init__.py b/dpctl/program/__init__.py index e1ef0c5809..bfd6c2aa1d 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -27,33 +27,23 @@ import warnings from dpctl.compiler import ( - SpecializationConstant, 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. " 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/program/utils/__init__.py b/dpctl/program/utils/__init__.py deleted file mode 100644 index fdf5f8e67f..0000000000 --- a/dpctl/program/utils/__init__.py +++ /dev/null @@ -1,36 +0,0 @@ -# Data Parallel Control (dpctl) -# -# Copyright 2026 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. - -""" -A collection of utility functions for dpctl.program module (deprecated, use -dpctl.compiler.utils instead). -""" - -import warnings - -from dpctl.compiler.utils import SpecializationConstantInfo, parse_spirv_specializations - -__all__ = [ - "parse_spirv_specializations", - "SpecializationConstantInfo", -] - -warnings.warn( - "dpctl.program.utils is deprecated and will be removed in a future release." - " Use dpctl.compiler.utils instead.", - DeprecationWarning, - stacklevel=2, -) diff --git a/setup.py b/setup.py index 169ccfc8f0..e8f14e1204 100644 --- a/setup.py +++ b/setup.py @@ -29,7 +29,6 @@ "dpctl.compiler.utils", "dpctl.memory", "dpctl.program", - "dpctl.program.utils", "dpctl.utils", ], package_data={ From 08b61f0083618985ea3bf2b5ba3904823fae78fc Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:39:05 -0700 Subject: [PATCH 05/12] _program.pyx to _compiler.pyx also removes attempted header installation from program submodule in setup.py --- .flake8 | 2 +- dpctl/compiler/__init__.pxd | 4 ++-- dpctl/compiler/__init__.py | 2 +- dpctl/compiler/{_program.pxd => _compiler.pxd} | 0 dpctl/compiler/{_program.pyx => _compiler.pyx} | 0 dpctl/program/__init__.pxd | 11 +++++++++-- setup.py | 3 +-- 7 files changed, 14 insertions(+), 8 deletions(-) rename dpctl/compiler/{_program.pxd => _compiler.pxd} (100%) rename dpctl/compiler/{_program.pyx => _compiler.pyx} (100%) diff --git a/.flake8 b/.flake8 index 278bc6ae29..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/compiler/_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/dpctl/compiler/__init__.pxd b/dpctl/compiler/__init__.pxd index f9da512601..b3d6afb379 100644 --- a/dpctl/compiler/__init__.pxd +++ b/dpctl/compiler/__init__.pxd @@ -18,11 +18,11 @@ # cython: language_level=3 """Declares the extension types and functions for the Cython API -implemented in dpctl.compiler._program.pyx. +implemented in dpctl.compiler._compiler.pyx. """ -from dpctl.compiler._program cimport ( +from dpctl.compiler._compiler cimport ( SyclKernel, SyclKernelBundle, create_kernel_bundle_from_source, diff --git a/dpctl/compiler/__init__.py b/dpctl/compiler/__init__.py index d6d3f8f0d2..3c16395140 100644 --- a/dpctl/compiler/__init__.py +++ b/dpctl/compiler/__init__.py @@ -21,7 +21,7 @@ """ -from dpctl.compiler._program import ( +from dpctl.compiler._compiler import ( SpecializationConstant, SyclKernel, SyclKernelBundle, diff --git a/dpctl/compiler/_program.pxd b/dpctl/compiler/_compiler.pxd similarity index 100% rename from dpctl/compiler/_program.pxd rename to dpctl/compiler/_compiler.pxd diff --git a/dpctl/compiler/_program.pyx b/dpctl/compiler/_compiler.pyx similarity index 100% rename from dpctl/compiler/_program.pyx rename to dpctl/compiler/_compiler.pyx diff --git a/dpctl/program/__init__.pxd b/dpctl/program/__init__.pxd index 688b002467..25f1e0580c 100644 --- a/dpctl/program/__init__.pxd +++ b/dpctl/program/__init__.pxd @@ -19,9 +19,16 @@ """ Declares the extension types and functions for the Cython API -implemented in dpctl.compiler._program.pyx (deprecated, use dpctl.compiler +implemented in dpctl.compiler._compiler.pyx (deprecated, use dpctl.compiler instead). """ -from dpctl.compiler._program cimport * +from dpctl.compiler._compiler cimport ( + SyclKernel, + SyclKernelBundle, + create_kernel_bundle_from_source, + create_kernel_bundle_from_spirv, + create_program_from_source, + create_program_from_spirv, +) diff --git a/setup.py b/setup.py index e8f14e1204..4042ec0075 100644 --- a/setup.py +++ b/setup.py @@ -45,8 +45,7 @@ "include/syclinterface/Support/*.h", "include/dpctl/_sycl*.h", "include/dpctl/memory/_memory*.h", - "include/dpctl/compiler/_program*.h", - "include/dpctl/program/_program*.h", + "include/dpctl/compiler/_compiler*.h", "*.pxd", "compiler/*.pxd", "memory/*.pxd", From 8de00391a898a64d42beb99be11eb4eaca5c4993 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:49:32 -0700 Subject: [PATCH 06/12] remove program from CMake no longer contains sources --- dpctl/CMakeLists.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/dpctl/CMakeLists.txt b/dpctl/CMakeLists.txt index 3b259927b9..a4378a03d0 100644 --- a/dpctl/CMakeLists.txt +++ b/dpctl/CMakeLists.txt @@ -204,6 +204,5 @@ target_include_directories(${_trgt} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}) target_link_libraries(DpctlCAPI INTERFACE ${_trgt}_headers) add_subdirectory(compiler) -add_subdirectory(program) add_subdirectory(memory) add_subdirectory(utils) From 6a347dff29fd940a4dc1bdf4a6010d29243add47 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:49:36 -0700 Subject: [PATCH 07/12] remove remaining references to dpctl.program --- dpctl/_sycl_queue.pxd | 2 +- dpctl/_sycl_queue.pyx | 10 +++++----- dpctl/apis/include/dpctl_capi.h | 6 +++--- dpctl/compiler/CMakeLists.txt | 1 - dpctl/compiler/_compiler.pxd | 4 ---- dpctl/compiler/_compiler.pyx | 10 +++++----- dpctl/tests/test_raw_kernel_arg.py | 2 +- dpctl/tests/test_work_group_memory.py | 2 +- dpctl/tests/test_work_group_memory_opencl.py | 2 +- examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt | 2 +- 10 files changed, 18 insertions(+), 23 deletions(-) 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/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/compiler/CMakeLists.txt b/dpctl/compiler/CMakeLists.txt index e6feca320c..da92df546d 100644 --- a/dpctl/compiler/CMakeLists.txt +++ b/dpctl/compiler/CMakeLists.txt @@ -1,4 +1,3 @@ - file(GLOB _cython_sources *.pyx) foreach(_cy_file ${_cython_sources}) get_filename_component(_trgt ${_cy_file} NAME_WLE) diff --git a/dpctl/compiler/_compiler.pxd b/dpctl/compiler/_compiler.pxd index de75753495..2f63c308fb 100644 --- a/dpctl/compiler/_compiler.pxd +++ b/dpctl/compiler/_compiler.pxd @@ -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/compiler/_compiler.pyx b/dpctl/compiler/_compiler.pyx index 1f3d154cca..623450fd24 100644 --- a/dpctl/compiler/_compiler.pyx +++ b/dpctl/compiler/_compiler.pyx @@ -187,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) @@ -253,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. """ @@ -273,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`: 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_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}) From 04ba2ad529d2ed650d3ddc0fe8223c19fc839805 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 19:17:37 -0700 Subject: [PATCH 08/12] remove Cython API from program --- .../api_reference/dpctl_cython.rst | 10 +++--- dpctl/compiler/__init__.pxd | 2 -- dpctl/program/__init__.pxd | 34 ------------------- setup.py | 1 - 4 files changed, 5 insertions(+), 42 deletions(-) delete mode 100644 dpctl/program/__init__.pxd 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/dpctl/compiler/__init__.pxd b/dpctl/compiler/__init__.pxd index b3d6afb379..ce78a4af38 100644 --- a/dpctl/compiler/__init__.pxd +++ b/dpctl/compiler/__init__.pxd @@ -27,6 +27,4 @@ from dpctl.compiler._compiler cimport ( SyclKernelBundle, create_kernel_bundle_from_source, create_kernel_bundle_from_spirv, - create_program_from_source, - create_program_from_spirv, ) diff --git a/dpctl/program/__init__.pxd b/dpctl/program/__init__.pxd deleted file mode 100644 index 25f1e0580c..0000000000 --- a/dpctl/program/__init__.pxd +++ /dev/null @@ -1,34 +0,0 @@ -# Data Parallel Control (dpctl) -# -# Copyright 2020 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. - -# distutils: language = c++ -# cython: language_level=3 - -""" -Declares the extension types and functions for the Cython API -implemented in dpctl.compiler._compiler.pyx (deprecated, use dpctl.compiler -instead). -""" - - -from dpctl.compiler._compiler cimport ( - SyclKernel, - SyclKernelBundle, - create_kernel_bundle_from_source, - create_kernel_bundle_from_spirv, - create_program_from_source, - create_program_from_spirv, -) diff --git a/setup.py b/setup.py index 4042ec0075..d4c9987b1d 100644 --- a/setup.py +++ b/setup.py @@ -49,7 +49,6 @@ "*.pxd", "compiler/*.pxd", "memory/*.pxd", - "program/*.pxd", ] }, include_package_data=False, From 28ce170dd87b0f9c49635e22ac5dd0beb348d291 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 19:17:49 -0700 Subject: [PATCH 09/12] apply review feedback --- docs/doc_sources/api_reference/dpctl/compiler.rst | 2 +- docs/doc_sources/api_reference/dpctl/program.rst | 2 +- dpctl/compiler/_compiler.pxd | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/docs/doc_sources/api_reference/dpctl/compiler.rst b/docs/doc_sources/api_reference/dpctl/compiler.rst index c887ecf562..fae4c27373 100644 --- a/docs/doc_sources/api_reference/dpctl/compiler.rst +++ b/docs/doc_sources/api_reference/dpctl/compiler.rst @@ -8,7 +8,7 @@ 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`. diff --git a/docs/doc_sources/api_reference/dpctl/program.rst b/docs/doc_sources/api_reference/dpctl/program.rst index 619947fd6c..61dc1839d4 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -11,7 +11,7 @@ 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`. diff --git a/dpctl/compiler/_compiler.pxd b/dpctl/compiler/_compiler.pxd index 2f63c308fb..47af0a604a 100644 --- a/dpctl/compiler/_compiler.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. """ From 5e670eb974cc82f4887dd1183751713831c64e64 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 19:21:19 -0700 Subject: [PATCH 10/12] fix namespace in test_sycl_program --- dpctl/tests/test_sycl_program.py | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 4cc8d8a4fc..cc7292a021 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -343,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 @@ -402,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 From cbc558833dd3d1fd9253c2238fd49428cb6cf392 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 5 Aug 2026 19:21:25 -0700 Subject: [PATCH 11/12] add gh-2317 to changelog --- CHANGELOG.md | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index c033895c19..e652b7ba97 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,9 @@ 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) + ### Maintenance * Updated pybind11 version used by `dpctl` and examples [gh-2357](https://github.com/IntelPython/dpctl/pull/2357) From 30bba23987e62d6c2f23f73de2bb83f1ad29fb10 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Mon, 10 Aug 2026 14:09:29 -0700 Subject: [PATCH 12/12] remove and deprecate remaining references to program --- CHANGELOG.md | 1 + .../contributor_guides/building.rst | 2 +- dpctl/apis/include/dpctl4pybind11.hpp | 2 +- ..._sycl_program.py => test_sycl_compiler.py} | 48 +++++++++---------- .../pybind11/use_dpctl_sycl_kernel/example.py | 10 ++-- libsyclinterface/CMakeLists.txt | 28 ++++++++--- libsyclinterface/dbg_build.sh | 2 +- libsyclinterface/dbg_build_custom.sh | 2 +- .../syclinterface/Config/dpctl_config.h.in | 5 +- .../dpctl_sycl_kernel_bundle_interface.cpp | 26 +++++----- .../test_sycl_kernel_bundle_interface.cpp | 2 +- scripts/_build_helper.py | 3 +- 12 files changed, 74 insertions(+), 57 deletions(-) rename dpctl/tests/{test_sycl_program.py => test_sycl_compiler.py} (92%) diff --git a/CHANGELOG.md b/CHANGELOG.md index e652b7ba97..9ac87f4eb1 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -22,6 +22,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### 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/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/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index f23f090de8..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 *); diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_compiler.py similarity index 92% rename from dpctl/tests/test_sycl_program.py rename to dpctl/tests/test_sycl_compiler.py index cc7292a021..d7364d8de1 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_compiler.py @@ -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 dpc.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,41 +54,41 @@ 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 dpc.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 @@ -114,7 +114,7 @@ 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 @@ -149,7 +149,7 @@ def _check_cpython_api_SyclKernel_Make(krn): assert krn.work_group_size == k3.work_group_size -def _check_multi_kernel_program(kb): +def _check_multi_kernel_bundle(kb): assert type(kb) is dpc.SyclKernelBundle assert type(kb.addressof_ref()) is int @@ -206,7 +206,7 @@ def test_create_kernel_bundle_from_source_ocl(): except dpctl.SyclQueueCreationError: pytest.skip("No OpenCL queue is available") kb = dpc.create_kernel_bundle_from_source(q, oclSrc) - _check_multi_kernel_program(kb) + _check_multi_kernel_bundle(kb) def test_create_kernel_bundle_from_spirv_ocl(): @@ -218,7 +218,7 @@ def test_create_kernel_bundle_from_spirv_ocl(): with open(spirv_file, "rb") as fin: spirv = fin.read() kb = dpc.create_kernel_bundle_from_spirv(q, spirv) - _check_multi_kernel_program(kb) + _check_multi_kernel_bundle(kb) def test_create_kernel_bundle_from_spirv_l0(): @@ -230,7 +230,7 @@ def test_create_kernel_bundle_from_spirv_l0(): with open(spirv_file, "rb") as fin: spirv = fin.read() kb = dpc.create_kernel_bundle_from_spirv(q, spirv) - _check_multi_kernel_program(kb) + _check_multi_kernel_bundle(kb) @pytest.mark.xfail( @@ -251,7 +251,7 @@ def test_create_kernel_bundle_from_source_l0(): c[index] = a[index] + d*b[index]; \ }" kb = dpc.create_kernel_bundle_from_source(q, oclSrc) - _check_multi_kernel_program(kb) + _check_multi_kernel_bundle(kb) def test_create_kernel_bundle_from_invalid_src_ocl(): diff --git a/examples/pybind11/use_dpctl_sycl_kernel/example.py b/examples/pybind11/use_dpctl_sycl_kernel/example.py index 10d6b39e12..e0f7276954 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/example.py +++ b/examples/pybind11/use_dpctl_sycl_kernel/example.py @@ -30,12 +30,12 @@ with open("resource/double_it.spv", "br") as fh: il = fh.read() -# Build the program for the selected device -pr = dpc.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 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'}", ]