From 794615212e137ce514013ebad22fa217ea2dd7d1 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Mon, 27 Apr 2026 19:15:23 -0700 Subject: [PATCH 01/14] Implement SpecializationConstant class --- dpctl/_backend.pxd | 4 + dpctl/_sycl_platform.pyx | 2 +- dpctl/program/__init__.py | 2 + dpctl/program/_program.pyx | 170 ++++++++++++++++++ .../dpctl_sycl_kernel_bundle_interface.h | 7 + 5 files changed, 184 insertions(+), 1 deletion(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 93d9b5ef97..d2f1a0decb 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -431,6 +431,10 @@ cdef extern from "syclinterface/dpctl_sycl_context_interface.h": cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": + ctypedef struct _spec_const "DPCTLSpecConst": + uint32_t id + size_t size + const void *value cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSpirv( const DPCTLSyclContextRef Ctx, const DPCTLSyclDeviceRef Dev, diff --git a/dpctl/_sycl_platform.pyx b/dpctl/_sycl_platform.pyx index 41eff7b5d3..ba78226e50 100644 --- a/dpctl/_sycl_platform.pyx +++ b/dpctl/_sycl_platform.pyx @@ -236,7 +236,7 @@ cdef class SyclPlatform(_SyclPlatform): and filter string for each device is printed. Args: - verbosity (Literal[0, 1, 2], optional):. + verbosity (Literal[0, 1, 2], optional): The verbosity controls how much information is printed by the function. Value ``0`` is the lowest level set by default and ``2`` is the highest level to print the most verbose output. diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index 71302e4186..bb16f9611b 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -22,6 +22,7 @@ """ from ._program import ( + SpecializationConstant, SyclKernel, SyclKernelBundle, SyclKernelBundleCompilationError, @@ -41,6 +42,7 @@ "SyclKernelBundleCompilationError", "SyclProgram", "SyclProgramCompilationError", + "SpecializationConstant", ] diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 8737be4762..7cecc3fa4d 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -26,7 +26,17 @@ an OpenCL source string or a SPIR-V binary file. """ +from cpython.buffer cimport ( + Py_buffer, + PyBUF_ANY_CONTIGUOUS, + PyBUF_SIMPLE, + PyBuffer_Release, + PyObject_CheckBuffer, + PyObject_GetBuffer, +) +from cpython.bytes cimport PyBytes_FromStringAndSize from libc.stdint cimport uint32_t +from libc.string cimport memcmp import warnings @@ -51,14 +61,20 @@ from dpctl._backend cimport ( # noqa: E211, E402; DPCTLSyclDeviceRef, DPCTLSyclKernelBundleRef, DPCTLSyclKernelRef, + _spec_const, ) +import numbers + +import numpy as np + __all__ = [ "create_kernel_bundle_from_source", "create_kernel_bundle_from_spirv", "SyclKernel", "SyclKernelBundle", "SyclKernelBundleCompilationError", + "SpecializationConstant", ] cdef class SyclKernelBundleCompilationError(Exception): @@ -252,6 +268,160 @@ cdef api SyclKernelBundle SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef): return SyclKernelBundle._create(copied_KBRef) +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. + + There are multiple ways to create a :class:`.SpecializationConstant`: + + - ``SpecializationConstant(spec_id, obj)`` + If the constructor is invoked with a single variadic argument, the + argument is expected to either expose the Python buffer protocol or be + coercible to a NumPy array. If the argument is coercible to a NumPy array + or is one, it must have a supported data type (bool, integral, floating + point, or void). The specialization constant will be constructed from the + data in the buffer + + - ``SpecializationConstant(spec_id, dtype, obj)`` + If the constructor is invoked with two variadic arguments, and the first + argument is a string, it is interpreted as a NumPy ``dtype`` string and the + second argument will be coerced to a NumPy array with that data type. + The data type specified by the first argument must be a supported data + type (bool, integral, floating point, or void). + + - ``SpecializationConstant(spec_id, nbytes, raw_ptr)`` + If the constructor is invoked with two variadic arguments where both are + integers, the first argument is interpreted as the number of bytes and + the second argument is interpreted as a pointer to the data. + + Note that when constructing from a buffer, the + :class:`.SpecializationConstant`, shares memory with the original object. + Modifications to the original object's data after construction will be + reflected when the :class:`.SpecializationConstant` is used to create a + :class:`.SyclKernelBundle`. This is not the case when constructing from a + raw pointer, as the data is copied. + + Args: + spec_id (int): + The SPIR-V specialization ID. + args: + Variadic argument, see class documentation. + + Raises: + TypeError: In case of incorrect arguments given to constructor, + failure to coerce to a buffer, or unsupported data type when + coercing to a buffer. + ValueError: If the provided object fails to construct a buffer. + """ + + cdef _spec_const _spec_const + cdef Py_buffer _buffer + + def __cinit__(self, spec_id, *args): + cdef int ret_code = 0 + cdef object target_obj = None + + if not isinstance(spec_id, numbers.Integral): + raise TypeError( + "Specialization constant ID must be of type `int`, got " + f"{type(spec_id)}" + ) + + if len(args) == 0 or len(args) > 2: + raise TypeError( + f"Constructor takes 2 or 3 arguments, got {len(args)}." + ) + + self._spec_const.id = spec_id + + if len(args) == 2: + if ( + isinstance(args[0], numbers.Integral) and + isinstance(args[1], numbers.Integral) + ): + target_obj = PyBytes_FromStringAndSize( + args[1], args[0] + ) + elif isinstance(args[0], str): + target_obj = np.ascontiguousarray(args[1], dtype=args[0]) + + elif len(args) == 1: + target_obj = args[0] + if not PyObject_CheckBuffer(target_obj): + # attempt to coerce to a numpy array + target_obj = np.ascontiguousarray(target_obj) + else: + raise TypeError( + "Invalid arguments." + ) + + if isinstance(target_obj, np.ndarray): + if target_obj.dtype.kind not in ("b", "i", "u", "f", "c", "V"): + raise TypeError( + "Coercion of input to buffer resulted in an unsupported " + f"data type '{target_obj.dtype}'. When coercing objects, " + "`SpecializationConstant` expects the data to coerce to a " + "supported type: bool, integral, real or complex floating " + "point, or void. To pass arbitrary data, use a " + "`memoryview` or `bytes` object, or pass the pointer and " + "size directly." + ) + + ret_code = PyObject_GetBuffer( + target_obj, &(self._buffer), PyBUF_SIMPLE | PyBUF_ANY_CONTIGUOUS + ) + if ret_code != 0: + raise ValueError( + "Failed to get buffer view for the provided object." + ) + self._spec_const.value = self._buffer.buf + self._spec_const.size = self._buffer.len + + def __dealloc__(self): + PyBuffer_Release(&(self._buffer)) + + def __repr__(self): + return f"SpecializationConstant({self._spec_const.id})" + + def __eq__(self, other): + if not isinstance(other, SpecializationConstant): + return False + cdef SpecializationConstant _other = other + if ( + self._spec_const.id != _other._spec_const.id or + self._spec_const.size != _other._spec_const.size or + self._spec_const.value != _other._spec_const.value + ): + return False + return memcmp( + self._spec_const.value, + _other._spec_const.value, + self._spec_const.size + ) == 0 + + @property + def id(self): + """Returns the specialization ID for this specialization constant.""" + return self._spec_const.id + + @property + def size(self): + """ + Returns the size in bytes of the data for this specialization constant. + """ + return self._spec_const.size + + cdef size_t addressof(self): + """ + Returns the address of the _spec_const for this + :class:`.SpecializationConstant` cast to ``size_t``. + """ + return &(self._spec_const) + + cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""): """ Creates a Sycl interoperability kernel bundle from an OpenCL source diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h index 07a76c3fd8..de7543016e 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -35,6 +35,13 @@ DPCTL_C_EXTERN_C_BEGIN +typedef struct DPCTLSpecConstTy +{ + uint32_t id; + size_t size; + const void *value; +} DPCTLSpecConst; + /** * @defgroup KernelBundleInterface Kernel_bundle class C wrapper */ From b5befc3c1d73f63175be005c3c815de970b15a1d Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 28 Apr 2026 14:55:36 -0700 Subject: [PATCH 02/14] hook specialization constants into kernel bundle interface --- dpctl/_backend.pxd | 4 +- dpctl/program/_program.pxd | 5 +- dpctl/program/_program.pyx | 53 +++++++++++++-- .../dpctl_sycl_kernel_bundle_interface.h | 6 +- .../dpctl_sycl_kernel_bundle_interface.cpp | 66 +++++++++++++++++-- 5 files changed, 118 insertions(+), 16 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index d2f1a0decb..d23edaf506 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -440,7 +440,9 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": const DPCTLSyclDeviceRef Dev, const void *IL, size_t Length, - const char *CompileOpts) + const char *CompileOpts, + size_t NumSpecConsts, + const _spec_const *SpecConsts) cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource( const DPCTLSyclContextRef Ctx, const DPCTLSyclDeviceRef Dev, diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 435ef68521..41f781cecd 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -63,7 +63,10 @@ cpdef create_kernel_bundle_from_source ( SyclQueue q, unicode source, unicode copts=* ) cpdef create_kernel_bundle_from_spirv ( - SyclQueue q, const unsigned char[:] IL, unicode copts=* + SyclQueue q, + const unsigned char[:] IL, + unicode copts=*, + list specializations=*, ) cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) cpdef create_program_from_spirv ( diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 7cecc3fa4d..50ad8aac02 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -36,6 +36,7 @@ from cpython.buffer cimport ( ) from cpython.bytes cimport PyBytes_FromStringAndSize from libc.stdint cimport uint32_t +from libc.stdlib cimport free, malloc from libc.string cimport memcmp import warnings @@ -469,7 +470,10 @@ cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""): cpdef create_kernel_bundle_from_spirv( - SyclQueue q, const unsigned char[:] IL, str copts="" + SyclQueue q, + const unsigned char[:] IL, + str copts="", + list specializations=None, ): """ Creates a Sycl interoperability kernel bundle from an SPIR-V binary. @@ -487,7 +491,9 @@ cpdef create_kernel_bundle_from_spirv( copts (str, optional) Optional compilation flags that will be used when compiling the kernel bundle. Default: ``""``. - + specializations (list, optional) + A list of :class:`.SpecializationConstant` objects to be used + when creating the kernel bundle. Default: ``None``. Returns: kernel_bundle (:class:`.SyclKernelBundle`) A :class:`.SyclKernelBundle` object wrapping the @@ -506,11 +512,44 @@ cpdef create_kernel_bundle_from_spirv( cdef size_t length = IL.shape[0] cdef bytes bCOpts = copts.encode("utf8") cdef const char *COpts = bCOpts - KBref = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, dIL, length, COpts - ) - if KBref is NULL: - raise SyclKernelBundleCompilationError() + cdef size_t num_spconsts + cdef _spec_const *spconsts + cdef SpecializationConstant spconst + + if specializations is not None: + num_spconsts = len(specializations) + spconsts = <_spec_const *>( + malloc(num_spconsts * sizeof(_spec_const)) + ) + if spconsts == NULL: + raise MemoryError( + "Failed to allocate memory for specialization constants." + ) + for i, spconst in enumerate(specializations): + if not isinstance(spconst, SpecializationConstant): + free(spconsts) + raise TypeError( + "All items in specializations must be of type " + f"`SpecializationConstant`, got {type(spconst)}" + ) + spconsts[i] = spconst._spec_const + else: + num_spconsts = 0 + spconsts = NULL + try: + KBref = DPCTLKernelBundle_CreateFromSpirv( + CRef, + DRef, + dIL, + length, COpts, + num_spconsts, + spconsts, + ) + if KBref is NULL: + raise SyclKernelBundleCompilationError() + finally: + if spconsts != NULL: + free(spconsts) return SyclKernelBundle._create(KBref) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h index de7543016e..3909a1c3d1 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h @@ -58,6 +58,8 @@ typedef struct DPCTLSpecConstTy * @param Length The size of the IL binary in bytes. * @param CompileOpts Optional compiler flags used when compiling the * SPIR-V binary. + * @param NumSpecConsts The number of specialization constants. + * @param SpecConsts An array of specialization constants. * @return A new SyclKernelBundleRef pointer if the kernel_bundle creation * succeeded, else returns NULL. * @ingroup KernelBundleInterface @@ -68,7 +70,9 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef Ctx, __dpctl_keep const DPCTLSyclDeviceRef Dev, __dpctl_keep const void *IL, size_t Length, - const char *CompileOpts); + const char *CompileOpts, + size_t NumSpecConsts, + const DPCTLSpecConst *SpecConsts); /*! * @brief Create a Sycl kernel bundle from an OpenCL kernel source string. diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index 78c714ecbb..73e5d9ef87 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -31,6 +31,7 @@ #include "dpctl_error_handlers.h" #include "dpctl_sycl_type_casters.hpp" #include /* OpenCL headers */ +#include #include #include #include @@ -170,6 +171,21 @@ std::string _GetErrorCode_ocl_impl(cl_int code) } } +typedef cl_int (*clSetProgramSpecializationConstantFT)(cl_program, + cl_uint, + size_t, + const void *); +const char *clSetProgramSpecializationConstant_Name = + "clSetProgramSpecializationConstant"; +clSetProgramSpecializationConstantFT get_clSetProgramSpecializationConstant() +{ + static auto st_clSetProgramSpecializationConstantF = + cl_loader::get().getSymbol( + clSetProgramSpecializationConstant_Name); + + return st_clSetProgramSpecializationConstantF; +} + DPCTLSyclKernelBundleRef _CreateKernelBundle_common_ocl_impl(cl_program clProgram, const context &ctx, @@ -235,7 +251,9 @@ _CreateKernelBundleWithIL_ocl_impl(const context &ctx, const device &dev, const void *IL, size_t il_length, - const char *CompileOpts) + const char *CompileOpts, + size_t NumSpecConsts, + const DPCTLSpecConst *SpecConsts) { auto clCreateProgramWithILF = get_clCreateProgramWithIL(); if (clCreateProgramWithILF == nullptr) { @@ -257,6 +275,22 @@ _CreateKernelBundleWithIL_ocl_impl(const context &ctx, return nullptr; } + if (SpecConsts != nullptr && NumSpecConsts > 0) { + auto clSetProgramSpecConstF = get_clSetProgramSpecializationConstant(); + if (clSetProgramSpecConstF) { + for (size_t i = 0; i < NumSpecConsts; ++i) { + clSetProgramSpecConstF(clProgram, SpecConsts[i].id, + SpecConsts[i].size, SpecConsts[i].value); + } + } + else { + error_handler("clSetProgramSpecializationConstant is not available " + "in the OpenCL implementation.", + __FILE__, __func__, __LINE__); + return nullptr; + } + } + return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev, CompileOpts); } @@ -428,7 +462,9 @@ _CreateKernelBundleWithIL_ze_impl(const context &SyclCtx, const device &SyclDev, const void *IL, size_t il_length, - const char *CompileOpts) + const char *CompileOpts, + size_t NumSpecConsts, + const DPCTLSpecConst *SpecConsts) { auto zeModuleCreateFn = get_zeModuleCreate(); if (zeModuleCreateFn == nullptr) { @@ -444,8 +480,22 @@ _CreateKernelBundleWithIL_ze_impl(const context &SyclCtx, ZeDevice = get_native(SyclDev); // Specialization constants are not supported by DPCTL at the moment + std::vector spec_ids; + std::vector spec_values; + + if (SpecConsts != nullptr && NumSpecConsts > 0) { + spec_ids.reserve(NumSpecConsts); + spec_values.reserve(NumSpecConsts); + for (size_t i = 0; i < NumSpecConsts; ++i) { + spec_ids.push_back(SpecConsts[i].id); + spec_values.push_back(SpecConsts[i].value); + } + } ze_module_constants_t ZeSpecConstants = {}; - ZeSpecConstants.numConstants = 0; + ZeSpecConstants.numConstants = static_cast(NumSpecConsts); + ZeSpecConstants.pConstantIds = spec_ids.empty() ? nullptr : spec_ids.data(); + ZeSpecConstants.pConstantValues = + spec_values.empty() ? nullptr : spec_values.data(); // Populate the Level Zero module descriptions ze_module_desc_t ZeModuleDesc = {}; @@ -583,7 +633,9 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, __dpctl_keep const DPCTLSyclDeviceRef DevRef, __dpctl_keep const void *IL, size_t length, - const char *CompileOpts) + const char *CompileOpts, + size_t NumSpecConsts, + const DPCTLSpecConst *SpecConsts) { DPCTLSyclKernelBundleRef KBRef = nullptr; if (!CtxRef) { @@ -611,12 +663,14 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, switch (BE) { case backend::opencl: KBRef = _CreateKernelBundleWithIL_ocl_impl(*SyclCtx, *SyclDev, IL, - length, CompileOpts); + length, CompileOpts, + NumSpecConsts, SpecConsts); break; case backend::ext_oneapi_level_zero: #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION KBRef = _CreateKernelBundleWithIL_ze_impl(*SyclCtx, *SyclDev, IL, - length, CompileOpts); + length, CompileOpts, + NumSpecConsts, SpecConsts); break; #endif default: From 8ddc44025398b0feb484cae1ce4ceee2bf78852c Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Tue, 28 Apr 2026 15:29:29 -0700 Subject: [PATCH 03/14] add test for SpecializationConstant use in kernel --- .../specialization_constant_kernel.spv | Bin 0 -> 2288 bytes dpctl/tests/test_sycl_program.py | 38 ++++++++++++++++++ 2 files changed, 38 insertions(+) create mode 100644 dpctl/tests/input_files/specialization_constant_kernel.spv diff --git a/dpctl/tests/input_files/specialization_constant_kernel.spv b/dpctl/tests/input_files/specialization_constant_kernel.spv new file mode 100644 index 0000000000000000000000000000000000000000..d696fa755f7e56d61ceeab707c468f0c937b8ed6 GIT binary patch literal 2288 zcmZ{kTTc^F5XVoUih_cOco&O!Ll6NGwRj6qN~MBQgNZMjrQPTzWlP${fbqd6KZmcr z`Q=R11mpj=yT@!uJlaW#ZFsI!sm~{7}LVC%h-&*Aspt z;roNebek?ydAZ@NEU!D>?zUGA9ka^wQVn0C{MJ)n$ew|y1_)$XNJ^) zF(*aN+t-!Vnbuyle%INUnyq=E>({4e)*CxP&8w}u!8>ooYX)BZjUT>WSuU+sCY>2S z^tMW6f2vU4oGKIwNYj(^Zp*J$)X734XoYc7p`a_Gxwx-SoxvAH+}l1Uj_*{z;PUlL zM)^#n9H+JIH$OP@JAOSZ1;u(})2)|+4~?oD`i-EZF%@?-x}$Au#s3IChs+WEjH%WI z*#xua^BO*%Z!#|aZdGaHkI~8J-DcC>Bkq_v{x8mP-g$x7bVIM^w8O%fqN15&RMfV= z1GDfAi^74N(~+EwWcK1>q{n#vw_Y-@tpoP94l;EvM|HAhK>WIh9VrTV^l=e5p`Q{@ z`_a?>ymap2(DP&Qo&GqtAemUcC$*Li=w<2Y+%@Us4l0&@o{2co67u<0a+IXPCRcg`94JH@#lo|DeJ)I%STr%U{z ztl)0@9Fmt^!-mXk=;R+0CkGpPpE$b!21C->J!}yE6XL{UA6BFwjCAVYBj#LU<9=$e zN5~NVOX6{#v3^y0e815SbeP*o+`KrwfI38db~}12kYIE4Igpr$CwDLh|7syNkbkN z;p2QuX1pUA-0*RqtuZMXd=J)`k_{in)|i$|4Sc+p@4AogaaTlraQvl8{9k}M{D;00 zweN|D;XBNT&WP}le_uome6u2azl9-oo#dI~Xr*X;2wkE Date: Tue, 28 Apr 2026 16:25:17 -0700 Subject: [PATCH 04/14] fix libsyclinterface tests --- .../test_sycl_kernel_bundle_interface.cpp | 24 +++++++++++-------- .../tests/test_sycl_queue_submit.cpp | 6 +++-- ...t_sycl_queue_submit_local_accessor_arg.cpp | 6 +++-- .../test_sycl_queue_submit_raw_kernel_arg.cpp | 6 +++-- ...ycl_queue_submit_work_group_memory_arg.cpp | 6 +++-- 5 files changed, 30 insertions(+), 18 deletions(-) diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp index a835c277b9..5793e983d9 100644 --- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp @@ -69,7 +69,8 @@ struct TestDPCTLSyclKernelBundleInterface spirvFile.seekg(0, std::ios::beg); spirvFile.read(spirvBuffer.data(), spirvFileSize); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr); + CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr, 0, + nullptr); } } @@ -132,18 +133,21 @@ TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvNull) const void *null_spirv = nullptr; DPCTLSyclKernelBundleRef KBRef = nullptr; // Null context - EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( - Null_CRef, Null_DRef, null_spirv, 0, nullptr)); + EXPECT_NO_FATAL_FAILURE( + KBRef = DPCTLKernelBundle_CreateFromSpirv( + Null_CRef, Null_DRef, null_spirv, 0, nullptr, 0, nullptr)); ASSERT_TRUE(KBRef == nullptr); // Null device - EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, Null_DRef, null_spirv, 0, nullptr)); + EXPECT_NO_FATAL_FAILURE( + KBRef = DPCTLKernelBundle_CreateFromSpirv(CRef, Null_DRef, null_spirv, + 0, nullptr, 0, nullptr)); ASSERT_TRUE(KBRef == nullptr); // Null IL - EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, null_spirv, 0, nullptr)); + EXPECT_NO_FATAL_FAILURE( + KBRef = DPCTLKernelBundle_CreateFromSpirv(CRef, DRef, null_spirv, 0, + nullptr, 0, nullptr)); ASSERT_TRUE(KBRef == nullptr); } @@ -350,8 +354,8 @@ TEST_F(TestKernelBundleUnsupportedBackend, CheckCreateFromSpirv) spirvFile.close(); DPCTLSyclKernelBundleRef KBRef = nullptr; - EXPECT_NO_FATAL_FAILURE( - KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr)); + EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer.data(), spirvFileSize, + nullptr, 0, nullptr)); ASSERT_TRUE(KBRef == nullptr); } diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index ab5b6bef82..f2fc2b2140 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -242,7 +242,8 @@ struct TestQueueSubmit : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDevice_Delete(DRef); DPCTLDeviceSelector_Delete(DSRef); } @@ -282,7 +283,8 @@ struct TestQueueSubmitFP64 : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDeviceSelector_Delete(DSRef); } diff --git a/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp b/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp index f6110375bb..8f1b97d1d4 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp @@ -237,7 +237,8 @@ struct TestQueueSubmitWithLocalAccessor : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDevice_Delete(DRef); DPCTLDeviceSelector_Delete(DSRef); } @@ -276,7 +277,8 @@ struct TestQueueSubmitWithLocalAccessorFP64 : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDeviceSelector_Delete(DSRef); } diff --git a/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp b/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp index f40bc20066..04d3958d8d 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp @@ -262,7 +262,8 @@ struct TestQueueSubmitWithRawKernelArg : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDevice_Delete(DRef); DPCTLDeviceSelector_Delete(DSRef); } @@ -301,7 +302,8 @@ struct TestQueueSubmitWithRawKernelArgFP64 : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDeviceSelector_Delete(DSRef); } diff --git a/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp b/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp index d0f44b7275..d1d1f69bfa 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp @@ -262,7 +262,8 @@ struct TestQueueSubmitWithWorkGroupMemory : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDevice_Delete(DRef); DPCTLDeviceSelector_Delete(DSRef); } @@ -301,7 +302,8 @@ struct TestQueueSubmitWithWorkGroupMemoryFP64 : public ::testing::Test auto CRef = DPCTLQueue_GetContext(QRef); KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0, + nullptr); DPCTLDeviceSelector_Delete(DSRef); } From 1de2f077367439b32f8aff5181dc023e599be668 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Thu, 30 Apr 2026 11:35:16 -0700 Subject: [PATCH 05/14] add test for composite specialization constant also removes "v" as a permitted specialization constant intermediate data type, as composite specialization constants are broken into multiple specialization constants, so structs end up passed as a single constant while the program expects multiple, and therefore, doesn't work as intended --- dpctl/program/_program.pyx | 12 ++--- .../specialization_constant_composite.spv | Bin 0 -> 3432 bytes dpctl/tests/test_sycl_program.py | 41 ++++++++++++++++++ 3 files changed, 47 insertions(+), 6 deletions(-) create mode 100644 dpctl/tests/input_files/specialization_constant_composite.spv diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index 50ad8aac02..166d46a3e3 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -282,8 +282,8 @@ cdef class SpecializationConstant: If the constructor is invoked with a single variadic argument, the argument is expected to either expose the Python buffer protocol or be coercible to a NumPy array. If the argument is coercible to a NumPy array - or is one, it must have a supported data type (bool, integral, floating - point, or void). The specialization constant will be constructed from the + or is one, it must have a supported data type (bool, integral, or + floating point). The specialization constant will be constructed from the data in the buffer - ``SpecializationConstant(spec_id, dtype, obj)`` @@ -291,7 +291,7 @@ cdef class SpecializationConstant: argument is a string, it is interpreted as a NumPy ``dtype`` string and the second argument will be coerced to a NumPy array with that data type. The data type specified by the first argument must be a supported data - type (bool, integral, floating point, or void). + type (bool, integral, or floating point). - ``SpecializationConstant(spec_id, nbytes, raw_ptr)`` If the constructor is invoked with two variadic arguments where both are @@ -360,13 +360,13 @@ cdef class SpecializationConstant: ) if isinstance(target_obj, np.ndarray): - if target_obj.dtype.kind not in ("b", "i", "u", "f", "c", "V"): + if target_obj.dtype.kind not in ("b", "i", "u", "f", "c"): raise TypeError( "Coercion of input to buffer resulted in an unsupported " f"data type '{target_obj.dtype}'. When coercing objects, " "`SpecializationConstant` expects the data to coerce to a " - "supported type: bool, integral, real or complex floating " - "point, or void. To pass arbitrary data, use a " + "supported type: bool, integral, or real or complex " + "floating point. To pass arbitrary data, use a " "`memoryview` or `bytes` object, or pass the pointer and " "size directly." ) diff --git a/dpctl/tests/input_files/specialization_constant_composite.spv b/dpctl/tests/input_files/specialization_constant_composite.spv new file mode 100644 index 0000000000000000000000000000000000000000..6b69617a7dc599495bc6a43862832b1ef77eab0d GIT binary patch literal 3432 zcma)-ZBHCk6vt=FQYyt_ZE1aIk=oLNR4fnDx)numffb~brJzmPVb~pD?LOe{fPgir zFZ6TxR$uu={5U3RQse(Ob5Ge2ns}4{p2vI6x#ylc+j8P+o0m#?XS|c%4_>m)dMCWJ zSZJ%GCmlWQ=ygYLIQqOR@&Dt3=e2vMy!_&dpIcbY&ifnN{zlM<{7MjQ`ITBRl)Wvs z&&@9R;of#wi2R_@*r|kyby^uyqW8AL>eT#TBPx2Ttwoi1F*?0&RqA=K3;x69{K9Bs zzfd0WSBA!lVHA|g!=p>JooX>GE(O)ia4xJ@!}4k=+M1iqE#zJB28(Ql@IvMwINzW|f2Uldn^iL4BaSFWD3Ox2RtF_sh>I$(%@g{Kj^v zzUxo!l*&=AI#aH#2jyIKw^j(EQmvZPG3ED4lH)kveCdGxw5qqKL=~OGCF)F zQ|GwY+ZE*)d(xduCaBkgef&DSS6}$~{$^MW>p>J2{bPF(=Y3QDjkr^T%RALlb#pPO z2bC}i>kV&4(Ku(n@PE$A9nYvw?9s@GPscp-bt#sUe!n!H>xaeRZY&Nwao&mJq`h;Z z*9B&XJjgc%`iOj9wAqf_Y`^Se?iBx7QS!2obE3#eEotuy$6LNd$$dhH&Wr!*5qV8A z`7C~LL@peWzmeRW|A~_=f7!{FzahCf|3Q=7rL6c(2rT3?qV$G^{FZ1Bg4Td;3IBG% zYB1e39NZdk`+ASa1CrTi_K(FA&thc7Ydzy<*u|jVEM#hGU@(8IqW5D(tYGA*8%S>fNV6RIbTiwDIwzib(!#WK z-2c453=ofgGv{9l=*<^B?{z_-N2aeS68`s`?I*=VUKEIdp1r*7`vQHmy}!A8$vfMU zsO4W4KOvYecT3zA0sjXAK0msAtK!Lt-OlxucxG`@;QYw=Sxn}getj*_uax=)PrshY zAHRnJ`^N<9J$Hn@N3faWPEF4}H_rk0_gr71JA!>5eCb$^1-=Q7MAw@5Ks8OGEgs1N|PyYdNu{1IO literal 0 HcmV?d00001 diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 7b5f8db20e..61b9fee843 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -300,3 +300,44 @@ def test_create_kernel_bundle_with_spec_const(): ht_e.wait() assert np.all(y == 43) + + +def test_create_kernel_bundle_with_composite_spec_const(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Could not create default queue") + + # 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) + + 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( + q, spv_bytes, specializations=[sp1, sp2, sp3] + ) + kernel = kb.get_sycl_kernel("_ZTS21StructSpecConstKernel") + + n = 128 + x = np.ones(n, dtype="f4") + y = np.zeros_like(x) + + x_usm = dpctl.memory.MemoryUSMDevice(x.nbytes, queue=q) + y_usm = dpctl.memory.MemoryUSMDevice(y.nbytes, queue=q) + + e1 = q.memcpy_async(x_usm, x, x.nbytes) + e2 = q.submit(kernel, [x_usm, y_usm], [n], dEvents=[e1]) + e3 = q.memcpy_async(y, y_usm, y.nbytes, [e2]) + + ht_e = q._submit_keep_args_alive([x_usm], [e3]) + + e3.wait() + ht_e.wait() + + # 1.0 * 10 + 2.5 = 12.5 + assert np.all(y == 12.5) From b100cf50ab74f7f93899b652e83b0065bb6755b5 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Mon, 4 May 2026 20:27:13 -0700 Subject: [PATCH 06/14] add program.utils namespace with SPIRV parser --- dpctl/program/__init__.py | 5 ++ dpctl/program/utils/__init__.py | 25 ++++++++ dpctl/program/utils/_utils.py | 106 ++++++++++++++++++++++++++++++++ 3 files changed, 136 insertions(+) create mode 100644 dpctl/program/utils/__init__.py create mode 100644 dpctl/program/utils/_utils.py diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py index bb16f9611b..e1e625ff7a 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -45,6 +45,11 @@ "SpecializationConstant", ] +# add submodules +__all__ += [ + "utils", +] + def __getattr__(name): if name == "SyclProgram": diff --git a/dpctl/program/utils/__init__.py b/dpctl/program/utils/__init__.py new file mode 100644 index 0000000000..474f154f95 --- /dev/null +++ b/dpctl/program/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.program module. +""" + +from ._utils import parse_spirv_specializations + +__all__ = [ + "parse_spirv_specializations", +] diff --git a/dpctl/program/utils/_utils.py b/dpctl/program/utils/_utils.py new file mode 100644 index 0000000000..2a16802dea --- /dev/null +++ b/dpctl/program/utils/_utils.py @@ -0,0 +1,106 @@ +# 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. + +"""Implements various utilities for the dpctl.program module.""" + +from enum import IntEnum + +import numpy as np + + +class SpirvOpCode(IntEnum): + OpName = 5 + OpTypeBool = 20 + OpTypeInt = 21 + OpTypeFloat = 22 + OpSpecConstantTrue = 48 + OpSpecConstantFalse = 49 + OpSpecConstant = 50 + OpDecorate = 71 + + +class SpirvDecoration(IntEnum): + SpecId = 1 + + +def parse_spirv_specializations( + spv_bytes: bytes | bytearray | memoryview, +) -> dict[int, dict[str, str]]: + words = np.frombuffer(spv_bytes, dtype=np.uint32) + + # verify magic number + if len(words) < 5 or words[0] != 0x07230203: + raise ValueError("Invalid SPIR-V binary") + + types = {} + ids = {} + names = {} + constants = {} + + i = 5 # skip 5 word header + while i < len(words): + word = words[i] + opcode = word & 0xFFFF + word_count = word >> 16 + + if word_count == 0: + raise ValueError(f"Invalid SPIR-V instruction at word index {i}") + + if opcode == SpirvOpCode.OpTypeBool: + result_id = int(words[i + 1]) + types[result_id] = "?" + elif opcode == SpirvOpCode.OpTypeInt: + result_id = int(words[i + 1]) + width = int(words[i + 2]) + signed = int(words[i + 3]) + prefix = "i" if signed else "u" + types[result_id] = f"{prefix}{width // 8}" + elif opcode == SpirvOpCode.OpTypeFloat: + result_id = int(words[i + 1]) + width = int(words[i + 2]) + types[result_id] = f"f{width // 8}" + elif opcode in ( + SpirvOpCode.OpSpecConstant, + SpirvOpCode.OpSpecConstantTrue, + SpirvOpCode.OpSpecConstantFalse, + ): + type_id = int(words[i + 1]) + result_id = int(words[i + 2]) + constants[result_id] = type_id + elif opcode == SpirvOpCode.OpDecorate: + target_id = int(words[i + 1]) + decoration = int(words[i + 2]) + if decoration == SpirvDecoration.SpecId: + ids[target_id] = int(words[i + 3]) + elif opcode == SpirvOpCode.OpName: + target_id = int(words[i + 1]) + name_bytes = words[i + 2 : i + word_count].tobytes() + names[target_id] = name_bytes.split(b"\x00", 1)[0].decode("utf-8") + + i += word_count + + result = {} + for target_id, spec_id in ids.items(): + type_id = constants.get(target_id) + dtype_str = types.get(type_id, "unknown_type") + name = names.get(target_id, f"unnamed_spec_const_{spec_id}") + + result[spec_id] = { + "name": name, + "dtype": dtype_str, + } + + return result From e2e4826192d1450920184b79074b16f83c281e99 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Mon, 4 May 2026 23:51:48 -0700 Subject: [PATCH 07/14] Refactor specialiazation constants to use dataclass also adds spec_id, itemsize, and default_value fields --- dpctl/program/utils/_utils.py | 122 +++++++++++++++--- .../specialization_constant_composite.spv | Bin 3432 -> 3504 bytes dpctl/tests/test_sycl_program.py | 47 +++++++ setup.py | 1 + 4 files changed, 152 insertions(+), 18 deletions(-) diff --git a/dpctl/program/utils/_utils.py b/dpctl/program/utils/_utils.py index 2a16802dea..86df0855ad 100644 --- a/dpctl/program/utils/_utils.py +++ b/dpctl/program/utils/_utils.py @@ -16,6 +16,7 @@ """Implements various utilities for the dpctl.program module.""" +from dataclasses import dataclass from enum import IntEnum import numpy as np @@ -29,6 +30,7 @@ class SpirvOpCode(IntEnum): OpSpecConstantTrue = 48 OpSpecConstantFalse = 49 OpSpecConstant = 50 + OpFunction = 54 OpDecorate = 71 @@ -36,9 +38,52 @@ class SpirvDecoration(IntEnum): SpecId = 1 +@dataclass(frozen=True) +class SpecializationConstantInfo: + """Data class representing specialization constant information.""" + + spec_id: int + dtype: str + name: str + itemsize: int + default_value: int | float | bool | None + + def parse_spirv_specializations( spv_bytes: bytes | bytearray | memoryview, -) -> dict[int, dict[str, str]]: +) -> tuple[SpecializationConstantInfo]: + """ + Parses SPIR-V byte stream to extract information about specializations, + including the specialization IDs, types, names, and default values. + + Note that the dtype information may be imprecise, as the compiler may + choose to, for example, represent a bool as char, or may represent both + signed and unsigned integers as unsigned integer bit buckets of the same + length. + + Args: + spv_bytes (bytes | bytearray | memoryview): + the SPIR-V byte stream. + + Returns: + tuple[SpecializationConstantInfo]: + a tuple of parsed constants and their information represented by + `SpecializationConstantInfo` objects, sorted by their + specialization IDs. The length of the tuple is equal to the number + of specialization constants found. Each + `SpecializationConstantInfo` object contains the following + attributes: + + - `spec_id` (int): The specialization ID. + - `dtype` (str): A NumPy style string representing the data type. + - `itemsize` (int): The size of the specialization constant in + bytes. + - `name` (str): The variable name. If not preserved in the binary, + a default name in the format `unnamed_spec_const_{spec_id}` is + used. + - `default_value` (int | float | bool | None): The default value of + the specialization constant. If not specified, `None` is used. + """ words = np.frombuffer(spv_bytes, dtype=np.uint32) # verify magic number @@ -49,6 +94,7 @@ def parse_spirv_specializations( ids = {} names = {} constants = {} + defaults = {} i = 5 # skip 5 word header while i < len(words): @@ -59,27 +105,45 @@ def parse_spirv_specializations( if word_count == 0: raise ValueError(f"Invalid SPIR-V instruction at word index {i}") - if opcode == SpirvOpCode.OpTypeBool: + if opcode == SpirvOpCode.OpFunction: + # everything following is not relevant to specialization constant + # parsing, so we can stop parsing at this point + break + elif opcode == SpirvOpCode.OpTypeBool: result_id = int(words[i + 1]) - types[result_id] = "?" + types[result_id] = {"dtype": "?", "itemsize": 1} elif opcode == SpirvOpCode.OpTypeInt: result_id = int(words[i + 1]) width = int(words[i + 2]) signed = int(words[i + 3]) prefix = "i" if signed else "u" - types[result_id] = f"{prefix}{width // 8}" + types[result_id] = { + "dtype": f"{prefix}{width // 8}", + "itemsize": width // 8, + } elif opcode == SpirvOpCode.OpTypeFloat: result_id = int(words[i + 1]) width = int(words[i + 2]) - types[result_id] = f"f{width // 8}" - elif opcode in ( - SpirvOpCode.OpSpecConstant, - SpirvOpCode.OpSpecConstantTrue, - SpirvOpCode.OpSpecConstantFalse, - ): + types[result_id] = { + "dtype": f"f{width // 8}", + "itemsize": width // 8, + } + elif opcode == SpirvOpCode.OpSpecConstant: + type_id = int(words[i + 1]) + result_id = int(words[i + 2]) + constants[result_id] = type_id + literal_words = words[i + 3 : i + word_count] + defaults[result_id] = literal_words.tobytes() + elif opcode == SpirvOpCode.OpSpecConstantTrue: type_id = int(words[i + 1]) result_id = int(words[i + 2]) constants[result_id] = type_id + defaults[result_id] = True + elif opcode == SpirvOpCode.OpSpecConstantFalse: + type_id = int(words[i + 1]) + result_id = int(words[i + 2]) + constants[result_id] = type_id + defaults[result_id] = False elif opcode == SpirvOpCode.OpDecorate: target_id = int(words[i + 1]) decoration = int(words[i + 2]) @@ -92,15 +156,37 @@ def parse_spirv_specializations( i += word_count - result = {} + # a spec ID may appear multiple times in the same binary with different + # target IDs. We only need to keep one, so skip duplicates + unique_ids = set() + result = [] for target_id, spec_id in ids.items(): + if spec_id in unique_ids: + continue + unique_ids.add(spec_id) type_id = constants.get(target_id) - dtype_str = types.get(type_id, "unknown_type") + type_info = types.get(type_id, {"dtype": "unknown_type", "itemsize": 0}) name = names.get(target_id, f"unnamed_spec_const_{spec_id}") - result[spec_id] = { - "name": name, - "dtype": dtype_str, - } - - return result + dtype_str = type_info["dtype"] + raw_default = defaults.get(target_id) + default_value = None + if isinstance(raw_default, bytes): + try: + default_value = np.frombuffer(raw_default, dtype=dtype_str)[ + 0 + ].item() + except Exception: + default_value = None + + result.append( + SpecializationConstantInfo( + spec_id=spec_id, + dtype=dtype_str, + name=name, + itemsize=type_info["itemsize"], + default_value=default_value, + ) + ) + + return tuple(sorted(result, key=lambda x: x.spec_id)) diff --git a/dpctl/tests/input_files/specialization_constant_composite.spv b/dpctl/tests/input_files/specialization_constant_composite.spv index 6b69617a7dc599495bc6a43862832b1ef77eab0d..c262f97ff6dfe48767c2ce52a31aea146a0d6b3e 100644 GIT binary patch literal 3504 zcmZ{lYjaao6oz+5C|D3gZeFm4q7(!wy$Bct3pJG3T1rzuQRnnDIc=lKNlZ>FMNs&l zzr(Nk$v@)HaYmiN@qJGAqM1(FGjG;)uf6tKYoBCyoY>_u8Mn{%xhJl>4!9jICl_o4czV*)Ay1zj?)v|E#JPUA+m+|;2Bq1B;!LpI3YNo85;VeOHE1-ek?gxt`}EEE zAbQY>DoGG_I%|zcvG(Z1jOeXa6i?0!cao}8ZdqmGQtWpF%G9}6h2Z`|d3L`=N5|)zYjHKI&WG_zG#$0$sD7`OtWMu7&6dvxQdlp=>&;4-)S7Wg*OcE&=^5MlmTM2G&ra2Lw`irVk_~hBLP3`=6xv)D zd%vP=WA9d{Pzc-Ya09=+?&arxL9i0VQ9DecYOrlB(sN&zeipD)p3)|+u<*xL2N_Aq7Mn-%t<(aQzsobsi3*-5ESR8s|ao~ybb{Z$=4vL-> zm?6$VJ}Xd1Bu^dcC53>Txbkx|^86S)92g~Gh;)!MQRq>o*F{Z?8J>h56$FLqClP@)A5d-;7 zDsB2(&CH(~v5v~0oH-AgmrS3@2l=2ZyF`yk%Y5102V~{Vp@XMh$n4)M${uvc`$d^~ z@*0rL8$<`Nzaol1`s4CE;AP@qJEHh+isl9KcuSO6)SE9Ovr3QfL64od1ER#mj-2;0 zcOZ`wqVNFwJEHibM}AL~I}rDfyvDq~;Pn>wBgyRR&PmRF>}7Hx&NVN0?~rqglBowW z_YFnazvyvUd*}f^8NuW}FP~OD$YorhpTyw}vNE!9|IdUmMJ2~sfqB7)Il;!9ulgmedCbIz0`Uih4c(nty)007YR9Zz@v_wpJu^Geqvsoh{+j6Z9zC`B zM;Y?JLvNvfsiWlliNM+Pdr~+mU?Z;^0=Z!;3fTTte(2#}6Nm?IceyE^m{Y=khj%z+ z`bVxMfjhZOHszVUtEvHaofXI@BixXV{@)Uq0pd|_W`9mVk8e&u&l|p-@|*gmtoYA+ z+s}%LTo#Cdp1u1-7laAHV*c*?a@X7Lh}!x0#9#06rMJZWLcpK5iO*Ai-hJ_$iQVq? zrFdqMdvSkc{Oo+I*H_X}uZ-#iPraVVAHS~!_A^sfdwPU^P_UV!r>3XRy>meSo~h!O z1p7WLNyiEWz6pz>l^(t-o}NddeB<~&EDL-es8LpQMc65jXI`*(wJIJP%kry<=S^T^ zw(Tz8h{tAggluo)Tk)J}b*g)rxq`R3!mlNeht1WZ`0oVnM11RO+w1Yo3HI(f;yH`l zh6MbXw?x2Z{x)xG-e&W*E}7Wq?ah8K-eU4*iOC#35U^QHo5KxnV-D?oKNN2TdJ&KQWv+ A5C8xG literal 3432 zcma)-ZBHCk6vt=FQYyt_ZE1aIk=oLNR4fnDx)numffb~brJzmPVb~pD?LOe{fPgir zFZ6TxR$uu={5U3RQse(Ob5Ge2ns}4{p2vI6x#ylc+j8P+o0m#?XS|c%4_>m)dMCWJ zSZJ%GCmlWQ=ygYLIQqOR@&Dt3=e2vMy!_&dpIcbY&ifnN{zlM<{7MjQ`ITBRl)Wvs z&&@9R;of#wi2R_@*r|kyby^uyqW8AL>eT#TBPx2Ttwoi1F*?0&RqA=K3;x69{K9Bs zzfd0WSBA!lVHA|g!=p>JooX>GE(O)ia4xJ@!}4k=+M1iqE#zJB28(Ql@IvMwINzW|f2Uldn^iL4BaSFWD3Ox2RtF_sh>I$(%@g{Kj^v zzUxo!l*&=AI#aH#2jyIKw^j(EQmvZPG3ED4lH)kveCdGxw5qqKL=~OGCF)F zQ|GwY+ZE*)d(xduCaBkgef&DSS6}$~{$^MW>p>J2{bPF(=Y3QDjkr^T%RALlb#pPO z2bC}i>kV&4(Ku(n@PE$A9nYvw?9s@GPscp-bt#sUe!n!H>xaeRZY&Nwao&mJq`h;Z z*9B&XJjgc%`iOj9wAqf_Y`^Se?iBx7QS!2obE3#eEotuy$6LNd$$dhH&Wr!*5qV8A z`7C~LL@peWzmeRW|A~_=f7!{FzahCf|3Q=7rL6c(2rT3?qV$G^{FZ1Bg4Td;3IBG% zYB1e39NZdk`+ASa1CrTi_K(FA&thc7Ydzy<*u|jVEM#hGU@(8IqW5D(tYGA*8%S>fNV6RIbTiwDIwzib(!#WK z-2c453=ofgGv{9l=*<^B?{z_-N2aeS68`s`?I*=VUKEIdp1r*7`vQHmy}!A8$vfMU zsO4W4KOvYecT3zA0sjXAK0msAtK!Lt-OlxucxG`@;QYw=Sxn}getj*_uax=)PrshY zAHRnJ`^N<9J$Hn@N3faWPEF4}H_rk0_gr71JA!>5eCb$^1-=Q7MAw@5Ks8OGEgs1N|PyYdNu{1IO diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 61b9fee843..564f40bed9 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -23,6 +23,7 @@ import dpctl import dpctl.program as dpctl_prog +from dpctl.program.utils import parse_spirv_specializations def get_spirv_abspath(fn): @@ -341,3 +342,49 @@ def test_create_kernel_bundle_with_composite_spec_const(): # 1.0 * 10 + 2.5 = 12.5 assert np.all(y == 12.5) + + +def test_spirv_specializations_parser(): + spirv_file = get_spirv_abspath("specialization_constant_kernel.spv") + with open(spirv_file, "rb") as spv: + spv_bytes = spv.read() + spec_consts = parse_spirv_specializations(spv_bytes) + assert len(spec_consts) == 1 + assert spec_consts[0].dtype == "u4" + + spirv_file = get_spirv_abspath("specialization_constant_composite.spv") + with open(spirv_file, "rb") as spv: + spv_bytes = spv.read() + + spec_consts = parse_spirv_specializations(spv_bytes) + assert len(spec_consts) == 3 + spec_const0, spec_const1, spec_const2 = spec_consts + assert spec_const0.dtype == "u4" + assert spec_const0.itemsize == 4 + assert spec_const0.name == "unnamed_spec_const_0" + assert spec_const0.default_value == 1 + + assert spec_const1.dtype == "f4" + assert spec_const1.itemsize == 4 + assert spec_const1.name == "unnamed_spec_const_1" + assert spec_const1.default_value == 0 + + # compiler translates bool to char + assert spec_const2.dtype == "u1" + assert spec_const2.itemsize == 1 + assert spec_const2.name == "unnamed_spec_const_2" + assert spec_const2.default_value == 0 + + +def test_spirv_specializations_parser_no_spec_consts(): + spirv_file = get_spirv_abspath("multi_kernel.spv") + with open(spirv_file, "rb") as spv: + spv_bytes = spv.read() + spec_consts = parse_spirv_specializations(spv_bytes) + assert not spec_consts + + +def test_spirv_specializations_parser_invalid_spirv(): + invalid_spv = b"\x00\x01\x02\x03\x04\x05" + with pytest.raises(ValueError): + parse_spirv_specializations(invalid_spv) diff --git a/setup.py b/setup.py index 1b10322ef8..2c44bd4a11 100644 --- a/setup.py +++ b/setup.py @@ -27,6 +27,7 @@ "dpctl", "dpctl.memory", "dpctl.program", + "dpctl.program.utils", "dpctl.utils", ], package_data={ From 3ccd0e4662b4659962450da6cd7f5c6f15a439b8 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:16:38 -0700 Subject: [PATCH 08/14] 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 | 30 +++++--- 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 +- examples/python/usm_memory_allocation.py | 12 +-- examples/python/usm_memory_host_access.py | 8 +- examples/python/usm_memory_operation.py | 6 +- setup.py | 4 + 26 files changed, 243 insertions(+), 93 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 5e442a8fc4..b0860cd352 100644 --- a/dpctl/__init__.py +++ b/dpctl/__init__.py @@ -125,6 +125,7 @@ # add submodules __all__ += [ "memory", + "compiler", "program", "utils", ] diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp index b1c17e8e90..49b242d68d 100644 --- a/dpctl/apis/include/dpctl4pybind11.hpp +++ b/dpctl/apis/include/dpctl4pybind11.hpp @@ -227,7 +227,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_ = @@ -450,7 +450,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 @@ -469,7 +469,7 @@ template <> struct type_caster } else { throw py::type_error("Input is of unexpected type, expected " - "dpctl.program.SyclKernel"); + "dpctl.compiler.SyclKernel"); } } @@ -482,12 +482,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 <> @@ -511,7 +511,7 @@ struct type_caster> } else { throw py::type_error("Input is of unexpected type, expected " - "dpctl.program.SyclKernelBundle"); + "dpctl.compiler.SyclKernelBundle"); } } @@ -526,7 +526,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 86df0855ad..94b64834c5 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 4e734f18ed..2f5c3f1cc6 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 e1e625ff7a..d1258630ed 100644 --- a/dpctl/program/__init__.py +++ b/dpctl/program/__init__.py @@ -15,13 +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 ._program import ( +import warnings + +from dpctl.compiler import ( SpecializationConstant, SyclKernel, SyclKernelBundle, @@ -50,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 474f154f95..807b9e88da 100644 --- a/dpctl/program/utils/__init__.py +++ b/dpctl/program/utils/__init__.py @@ -15,11 +15,21 @@ # 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 parse_spirv_specializations +import warnings + +from dpctl.compiler.utils import parse_spirv_specializations __all__ = [ "parse_spirv_specializations", ] + +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 0cb222573b..1334c36b96 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 f23f88988b..9ac359a646 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 3ebab3244e..f18867b59c 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 564f40bed9..5b6a72672f 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 a164d2cd0b..e7678123bb 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 805670bd8f..94e49282b8 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 be771e1e31..cd5ddab5c3 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/examples/python/usm_memory_allocation.py b/examples/python/usm_memory_allocation.py index 0c23b8bd20..8259e7bef0 100644 --- a/examples/python/usm_memory_allocation.py +++ b/examples/python/usm_memory_allocation.py @@ -18,23 +18,23 @@ Demonstrates SYCL USM memory usage in Python using dpctl.memory. """ -import dpctl.memory as dpmem +import dpctl.memory as dpm # allocate USM-shared byte-buffer -ms = dpmem.MemoryUSMShared(16) +ms = dpm.MemoryUSMShared(16) # allocate USM-device byte-buffer -md = dpmem.MemoryUSMDevice(16) +md = dpm.MemoryUSMDevice(16) # allocate USM-host byte-buffer -mh = dpmem.MemoryUSMHost(16) +mh = dpm.MemoryUSMHost(16) # specify alignment -mda = dpmem.MemoryUSMDevice(128, alignment=16) +mda = dpm.MemoryUSMDevice(128, alignment=16) # allocate using given queue, # i.e. on the device and bound to the context stored in the queue -mdq = dpmem.MemoryUSMDevice(256, queue=mda.sycl_queue) +mdq = dpm.MemoryUSMDevice(256, queue=mda.sycl_queue) # information about device associate with USM buffer print("Allocation performed on device:") diff --git a/examples/python/usm_memory_host_access.py b/examples/python/usm_memory_host_access.py index f6af9c4305..6eb11c93d8 100644 --- a/examples/python/usm_memory_host_access.py +++ b/examples/python/usm_memory_host_access.py @@ -19,14 +19,14 @@ Python program. """ -import dpctl.memory as dpmem +import dpctl.memory as dpm # USM-shared and USM-host pointers are host-accessible, # meaning they are accessible from Python, therefore # they implement Python buffer protocol # allocate 1K of USM-shared buffer -ms = dpmem.MemoryUSMShared(1024) +ms = dpm.MemoryUSMShared(1024) # create memoryview into USM-shared buffer msv = memoryview(ms) @@ -36,7 +36,7 @@ ir = i % 256 msv[i] = ir**2 % 256 -mh = dpmem.MemoryUSMHost(64) +mh = dpm.MemoryUSMHost(64) mhv = memoryview(mh) # copy content of block of USM-shared buffer to @@ -47,7 +47,7 @@ print(list(mhv)) # USM-device buffer is not host accessible -md = dpmem.MemoryUSMDevice(16) +md = dpm.MemoryUSMDevice(16) try: mdv = memoryview(md) except Exception as e: diff --git a/examples/python/usm_memory_operation.py b/examples/python/usm_memory_operation.py index 7c52baca09..ac9dfb678c 100644 --- a/examples/python/usm_memory_operation.py +++ b/examples/python/usm_memory_operation.py @@ -20,10 +20,10 @@ import numpy as np -import dpctl.memory as dpmem +import dpctl.memory as dpm -ms = dpmem.MemoryUSMShared(32) -md = dpmem.MemoryUSMDevice(32) +ms = dpm.MemoryUSMShared(32) +md = dpm.MemoryUSMDevice(32) host_buf = np.random.randint(0, 42, dtype=np.uint8, size=32) diff --git a/setup.py b/setup.py index 2c44bd4a11..9399f3cb0e 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 2c8b5a32d3ddd66c8c58795f0e407053d249e741 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:18:03 -0700 Subject: [PATCH 09/14] 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 87fdfb16b2..52dc55219a 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 aee4b574a8..d5294d7e04 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 99f007da729a4a4cf1fc9a041e15891db61b0736 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:19:51 -0700 Subject: [PATCH 10/14] 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 d5ad0689b359cd90b571852484ed64387d244fb8 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:32:26 -0700 Subject: [PATCH 11/14] move deprecated APIs into program submodule --- .../api_reference/dpctl/compiler.rst | 3 +- .../api_reference/dpctl/program.rst | 16 +---- 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 | 35 ----------- setup.py | 1 - 8 files changed, 65 insertions(+), 127 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 d5294d7e04..619947fd6c 100644 --- a/docs/doc_sources/api_reference/dpctl/program.rst +++ b/docs/doc_sources/api_reference/dpctl/program.rst @@ -23,20 +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 - -.. autosummary:: - :toctree: generated - :nosignatures: - - SyclKernelBundleCompilationError + 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 166d46a3e3..ac5eaabcde 100644 --- a/dpctl/compiler/_program.pyx +++ b/dpctl/compiler/_program.pyx @@ -39,8 +39,6 @@ from libc.stdint cimport uint32_t from libc.stdlib cimport free, malloc from libc.string cimport memcmp -import warnings - from dpctl._backend cimport ( # noqa: E211, E402; DPCTLKernel_Copy, DPCTLKernel_Delete, @@ -552,33 +550,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 d1258630ed..9ab03b969b 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 807b9e88da..0000000000 --- a/dpctl/program/utils/__init__.py +++ /dev/null @@ -1,35 +0,0 @@ -# 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.program module (deprecated, use -dpctl.compiler.utils instead). -""" - -import warnings - -from dpctl.compiler.utils import parse_spirv_specializations - -__all__ = [ - "parse_spirv_specializations", -] - -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 9399f3cb0e..36ec25c9fe 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 cadb8e5334b2c0b8e617b34a5104f0d62aba2949 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:39:05 -0700 Subject: [PATCH 12/14] _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 2f5c3f1cc6..462d2f6e36 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 36ec25c9fe..0d9472aeff 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 af0d7f89a3b3af9f3622a3a37f891a90be06c5f9 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:49:32 -0700 Subject: [PATCH 13/14] 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 c82fa86905cd2030865b42be591b9eb9007d3183 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 23 May 2026 17:49:36 -0700 Subject: [PATCH 14/14] 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 9469415cb8..eb10df860a 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 a9d0a2c71c..eb6f8be8c4 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -383,7 +383,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`. """ @@ -1178,10 +1178,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 @@ -1339,10 +1339,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 9cf245aef5..b32c3ecf79 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 41f781cecd..04e329f159 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 ac5eaabcde..5a69720a1e 100644 --- a/dpctl/compiler/_compiler.pyx +++ b/dpctl/compiler/_compiler.pyx @@ -186,14 +186,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) @@ -252,14 +252,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. """ @@ -272,7 +272,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 335f268144..a9d7543fe7 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 5f33f74162..97c44ed73c 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 462c1940d1..178bc74e9d 100644 --- a/dpctl/tests/test_work_group_memory_opencl.py +++ b/dpctl/tests/test_work_group_memory_opencl.py @@ -45,7 +45,7 @@ def test_submit_work_group_memory_opencl(): except dpctl.SyclQueueCreationError: pytest.skip("OpenCL queue could not be created") - kb = dpctl.program.create_kernel_bundle_from_source(q, ocl_kernel_src) + kb = dpctl.compiler.create_kernel_bundle_from_source(q, ocl_kernel_src) kernel = kb.get_sycl_kernel("local_mem_kernel") local_size = 16 global_size = local_size * 8 diff --git a/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt index 5d8129581e..5bb2a98c6f 100644 --- a/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt +++ b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt @@ -1,7 +1,7 @@ cmake_minimum_required(VERSION 3.21...3.27 FATAL_ERROR) project(use_queue_device VERSION 0.1 LANGUAGES CXX - DESCRIPTION "Example of using dpctl.program.SyclKernel <-> sycl::kernel type casting") + DESCRIPTION "Example of using dpctl.compiler.SyclKernel <-> sycl::kernel type casting") set(DPCTL_CMAKE_MODULES_PATH "${CMAKE_SOURCE_DIR}/../../../cmake") set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${DPCTL_CMAKE_MODULES_PATH})