From b7f435b96fb7cc7530c85a1aa59208ed25c860fe Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sun, 8 Nov 2026 12:50:44 -0800 Subject: [PATCH 1/3] check kernel argument count only for interop kernels The number of arguments is only defined for kernels from OpenCL C source or SPIR-V bundles, so SyclKernel records the kind of bundle it comes from and submit checks the count only for those. DPCTLKernel_GetNumArgs no longer aborts when the query throws. --- CHANGELOG.md | 2 + dpctl/_sycl_queue.pyx | 34 ++++--- dpctl/compiler/_compiler.pxd | 17 +++- dpctl/compiler/_compiler.pyx | 39 +++++--- dpctl/tests/helper/__init__.py | 2 + dpctl/tests/helper/_helper.py | 21 ++++ dpctl/tests/test_sycl_kernel_submit.py | 96 +++++++++++++++++++ .../dpctl_sycl_kernel_interface.h | 3 +- .../source/dpctl_sycl_kernel_interface.cpp | 9 +- 9 files changed, 193 insertions(+), 30 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 6c7c188b78..050589214a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -14,6 +14,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Added `dpctl.SyclQueue.get_last_event` and `dpctl.SyclQueue.is_get_last_event_available`, and the `DPCTLQueue_GetLastEvent` and `DPCTLQueue_GetLastEvent_Available` C-API functions, wrapping the DPC++ `sycl_ext_oneapi_in_order_queue_events` extension [gh-2299](https://github.com/IntelPython/dpctl/pull/2299) ### Changed +* `dpctl.SyclQueue.submit_async` and `dpctl.SyclQueue.submit` check the number of arguments only for kernels from kernel bundles created from OpenCL C source or SPIR-V, the only kernels whose number of arguments SYCL defines; `dpctl.compiler.SyclKernel` records the kind of kernel bundle it comes from [gh-2425](https://github.com/IntelPython/dpctl/pull/2425) * Object lifetime during offload is now managed by a list of releases deferred until the events that gate them have completed, rather than by a `host_task`. `dpctl` keeps one thread for looking over the list and running the releases that have come due, started by the first release deferred and again in the child of a `fork`, which takes the GIL to do so and lets a release call into the interpreter and free USM allocations [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * `import_dpctl()` now returns 0 on success and -1 with a Python exception set on failure, and `dpctl4pybind11.hpp` raises that exception rather than failing on the first C-API call [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * The event returned by `dpctl.SyclQueue.keep_args_alive`, `dpctl::utils::keep_args_alive` and `dpctl.SyclQueue._submit_keep_args_alive` is now for an empty kernel that gates the deferred release rather than for a `host_task` that performs it, so its completion means that the objects are no longer in use rather than that they have been released [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) @@ -27,6 +28,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Removed the deprecated `dpctl.program` submodule, use `dpctl.compiler` instead [gh-2407](https://github.com/IntelPython/dpctl/pull/2407) ### Fixed +* `DPCTLKernel_GetNumArgs` returns -1 instead of aborting the process when the number of arguments can not be queried, and `dpctl.compiler.SyclKernel.get_num_args` raises `RuntimeError` [gh-2425](https://github.com/IntelPython/dpctl/pull/2425) * Removed stale `select_host_device` and `has_host_device` names from `dpctl.__all__`, which made `from dpctl import *` fail with `AttributeError` [gh-2410](https://github.com/IntelPython/dpctl/pull/2410) * A USM allocation that fails now drops the references held for offloaded tasks that have completed, and is attempted once more if that released anything, so that memory only waiting to be given up is not reported as unavailable [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * `dpctl.SyclPlatform.get_devices` no longer prints Unified Runtime warnings for `device_type="custom"` on non-OpenCL platforms, where custom devices are not supported [gh-2306](https://github.com/IntelPython/dpctl/pull/2306) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 07fcd0cf20..52076a7230 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -82,6 +82,7 @@ from ._backend cimport ( # noqa: E211 _md_local_accessor, _queue_property_type, ) +from .compiler._compiler cimport _INTEROP_ORIGIN from .memory._memory cimport _Memory import ctypes @@ -1540,8 +1541,13 @@ cdef class SyclQueue(_SyclQueue): Raises: ValueError: - If the number of arguments in ``args`` differs from the number - of arguments the kernel takes. + If ``kernel`` is from a kernel bundle created from OpenCL C + source or SPIR-V, and the number of arguments in ``args`` + differs from the number of arguments the kernel takes. The + arguments of other kernels are not checked: the backend may + reject a wrong number of arguments, raising + :class:`SyclKernelSubmitError`, or run the kernel with the + arguments that were not given left unset. SyclKernelSubmitError: If ``kernel`` does not reference a ``sycl::kernel``, or if submission to the queue fails. @@ -1564,13 +1570,14 @@ cdef class SyclQueue(_SyclQueue): "``sycl::kernel``." ) # the backend can not detect that the correct number of arguments were - # provided, so check it here - nKA = kernel.num_args - if len(args) != nKA: - raise ValueError( - f"Kernel '{kernel.get_function_name()}' takes {nKA} " - f"argument(s), got {len(args)}." - ) + # provided, so check it here for the kernels it defines that number for + if kernel._origin == _INTEROP_ORIGIN: + nKA = kernel.num_args + if len(args) != nKA: + raise ValueError( + f"Kernel '{kernel.get_function_name()}' takes {nKA} " + f"argument(s), got {len(args)}." + ) # Allocate the arrays to be sent to DPCTLQueue_Submit kargs = malloc(len(args) * sizeof(void*)) @@ -1718,8 +1725,13 @@ cdef class SyclQueue(_SyclQueue): Raises: ValueError: - If the number of arguments in ``args`` differs from the number - of arguments the kernel takes. + If ``kernel`` is from a kernel bundle created from OpenCL C + source or SPIR-V, and the number of arguments in ``args`` + differs from the number of arguments the kernel takes. The + arguments of other kernels are not checked: the backend may + reject a wrong number of arguments, raising + :class:`SyclKernelSubmitError`, or run the kernel with the + arguments that were not given left unset. SyclKernelSubmitError: If ``kernel`` does not reference a ``sycl::kernel``, or if submission to the queue fails. diff --git a/dpctl/compiler/_compiler.pxd b/dpctl/compiler/_compiler.pxd index 3958a2f83e..dda329559f 100644 --- a/dpctl/compiler/_compiler.pxd +++ b/dpctl/compiler/_compiler.pxd @@ -28,16 +28,26 @@ from .._sycl_device cimport SyclDevice from .._sycl_queue cimport SyclQueue +cdef enum _kernel_origin: + # kernel from a bundle created outside of dpctl.compiler + _UNKNOWN_ORIGIN = 0 + # kernel from an OpenCL C source or SPIR-V bundle + _INTEROP_ORIGIN = 1 + # kernel from a SYCL source bundle + _SYCL_SOURCE_ORIGIN = 2 + + cdef api class SyclKernel [object PySyclKernelObject, type PySyclKernelType]: ''' Wraps a sycl::kernel object created from an OpenCL interoperability kernel. ''' cdef DPCTLSyclKernelRef _kernel_ref cdef str _function_name + cdef int _origin cdef DPCTLSyclKernelRef get_kernel_ref (self) @staticmethod - cdef SyclKernel _create (DPCTLSyclKernelRef kref, str name) + cdef SyclKernel _create (DPCTLSyclKernelRef kref, str name, int origin=*) cdef api class SyclKernelBundle [ @@ -52,11 +62,10 @@ cdef api class SyclKernelBundle [ binary file. """ cdef DPCTLSyclKernelBundleRef _kernel_bundle_ref - cdef bint _is_sycl_source + cdef int _origin @staticmethod - cdef SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref, - bint _is_sycl_source) + cdef SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref, int origin) cdef DPCTLSyclKernelBundleRef get_kernel_bundle_ref (self) cpdef SyclKernel get_sycl_kernel(self, str kernel_name) diff --git a/dpctl/compiler/_compiler.pyx b/dpctl/compiler/_compiler.pyx index 58d7a56310..0ecb24f35e 100644 --- a/dpctl/compiler/_compiler.pyx +++ b/dpctl/compiler/_compiler.pyx @@ -121,10 +121,13 @@ cdef class SyclKernel: """ """ @staticmethod - cdef SyclKernel _create(DPCTLSyclKernelRef kref, str name): + cdef SyclKernel _create( + DPCTLSyclKernelRef kref, str name, int origin=_UNKNOWN_ORIGIN + ): cdef SyclKernel ret = SyclKernel.__new__(SyclKernel) ret._kernel_ref = kref ret._function_name = name + ret._origin = origin return ret def __dealloc__(self): @@ -137,8 +140,21 @@ cdef class SyclKernel: def get_num_args(self): """ Returns the number of arguments for this kernel function. + + The number of arguments is defined by the backend only for kernels + from a kernel bundle created from OpenCL C source or SPIR-V. + + Raises: + RuntimeError: + If the number of arguments can not be queried. """ - return DPCTLKernel_GetNumArgs(self._kernel_ref) + cdef size_t n = DPCTLKernel_GetNumArgs(self._kernel_ref) + if n == -1: + raise RuntimeError( + "The number of arguments of kernel " + f"'{self._function_name}' can not be queried." + ) + return n cdef DPCTLSyclKernelRef get_kernel_ref(self): """ Returns the ``DPCTLSyclKernelRef`` pointer for this SyclKernel. @@ -256,11 +272,10 @@ cdef class SyclKernelBundle: """ @staticmethod - cdef SyclKernelBundle _create(DPCTLSyclKernelBundleRef KBRef, - bint is_sycl_source): + cdef SyclKernelBundle _create(DPCTLSyclKernelBundleRef KBRef, int origin): cdef SyclKernelBundle ret = SyclKernelBundle.__new__(SyclKernelBundle) ret._kernel_bundle_ref = KBRef - ret._is_sycl_source = is_sycl_source + ret._origin = origin return ret def __dealloc__(self): @@ -293,7 +308,7 @@ cdef class SyclKernelBundle: cdef bytes name = kernel_name.encode("utf8") cdef DPCTLSyclKernelRef KRef = NULL - if self._is_sycl_source: + if self._origin == _SYCL_SOURCE_ORIGIN: KRef = DPCTLKernelBundle_GetSyclKernel( self._kernel_bundle_ref, name ) @@ -308,11 +323,11 @@ cdef class SyclKernelBundle: "name, or the module it was built from may require device " "libraries that were not linked into it." ) - return SyclKernel._create(KRef, kernel_name) + return SyclKernel._create(KRef, kernel_name, self._origin) def has_sycl_kernel(self, str kernel_name): name = kernel_name.encode("utf8") - if self._is_sycl_source: + if self._origin == _SYCL_SOURCE_ORIGIN: return DPCTLKernelBundle_HasSyclKernel( self._kernel_bundle_ref, name ) @@ -346,7 +361,7 @@ cdef api SyclKernelBundle SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef): reference. """ cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef) - return SyclKernelBundle._create(copied_KBRef, False) + return SyclKernelBundle._create(copied_KBRef, _UNKNOWN_ORIGIN) cdef class SpecializationConstant: @@ -597,7 +612,7 @@ cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""): if KBref is NULL: raise SyclKernelBundleCompilationError() - return SyclKernelBundle._create(KBref, False) + return SyclKernelBundle._create(KBref, _INTEROP_ORIGIN) cpdef create_kernel_bundle_from_spirv( @@ -683,7 +698,7 @@ cpdef create_kernel_bundle_from_spirv( if spconsts != NULL: free(spconsts) - return SyclKernelBundle._create(KBref, False) + return SyclKernelBundle._create(KBref, _INTEROP_ORIGIN) cpdef create_kernel_bundle_from_sycl_source(SyclQueue q, @@ -819,4 +834,4 @@ cpdef create_kernel_bundle_from_sycl_source(SyclQueue q, DPCTLVirtualHeaderList_Delete(VirtualHeaders) DPCTLKernelBuildLog_Delete(BuildLog) - return SyclKernelBundle._create(KBref, True) + return SyclKernelBundle._create(KBref, _SYCL_SOURCE_ORIGIN) diff --git a/dpctl/tests/helper/__init__.py b/dpctl/tests/helper/__init__.py index be5f8d2e33..470fa13882 100644 --- a/dpctl/tests/helper/__init__.py +++ b/dpctl/tests/helper/__init__.py @@ -22,6 +22,7 @@ has_cpu, has_gpu, has_sycl_platforms, + make_kernel_of_unknown_origin, ) __all__ = [ @@ -30,4 +31,5 @@ "has_gpu", "has_sycl_platforms", "get_queue_or_skip", + "make_kernel_of_unknown_origin", ] diff --git a/dpctl/tests/helper/_helper.py b/dpctl/tests/helper/_helper.py index 5857b23195..a3f9886a42 100644 --- a/dpctl/tests/helper/_helper.py +++ b/dpctl/tests/helper/_helper.py @@ -49,3 +49,24 @@ def get_queue_or_skip(args=()): except dpctl.SyclQueueCreationError: pytest.skip(f"Queue could not be created from {args}") return q + + +def make_kernel_of_unknown_origin(krn): + """Wraps the ``sycl::kernel`` of ``krn`` with the ``SyclKernel_Make`` + C-API function, as a kernel created outside of ``dpctl.compiler`` is. + """ + import ctypes + import sys + + mod = sys.modules[krn.__class__.__module__] + cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer + cap_ptr_fn.restype = ctypes.c_void_p + cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p] + make_fn_ptr = cap_ptr_fn( + mod.__pyx_capi__["SyclKernel_Make"], + b"struct PySyclKernelObject *(DPCTLSyclKernelRef, char const *)", + ) + make_fn = ctypes.PYFUNCTYPE( + ctypes.py_object, ctypes.c_void_p, ctypes.c_char_p + )(make_fn_ptr) + return make_fn(krn.addressof_ref(), krn.get_function_name().encode()) diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index 9c57186d81..81e700aaf2 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -27,6 +27,8 @@ import dpctl.memory as dpm from dpctl._sycl_queue import kernel_arg_type +from .helper import make_kernel_of_unknown_origin + @pytest.mark.parametrize( "ctype_str,dtype,ctypes_ctor", @@ -343,6 +345,82 @@ def test_submit_local_accessor_arg(): expected = np.arange(1, x.size + 1, dtype=x.dtype) * (2 * lws) assert np.all(res == expected) + assert krn.num_args == 2 + with pytest.raises(ValueError): + q.submit_async(krn, [x_usm], [gws], [lws]) + q.submit( + make_kernel_of_unknown_origin(krn), + [x_usm, dpctl.LocalAccessor("i8", (lws,))], + [gws], + [lws], + ).wait() + + +_local_accessor_sycl_source = """ +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +extern "C" SYCL_EXTERNAL +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void reverse_in_group(sycl::local_accessor lmem, int *x, int scale) { + auto item = syclext::this_work_item::get_nd_item<1>(); + size_t lid = item.get_local_linear_id(); + size_t lws = item.get_local_range(0); + lmem[lid] = x[item.get_global_linear_id()]; + sycl::group_barrier(item.get_group()); + x[item.get_global_linear_id()] = scale * lmem[lws - 1 - lid]; +} +""" + + +@pytest.mark.parametrize("backend", ["opencl", "level_zero"]) +def test_submit_local_accessor_arg_sycl_source(backend): + try: + q = dpctl.SyclQueue(backend) + except dpctl.SyclQueueCreationError: + pytest.skip(f"No {backend} queue is available") + if not dpc.is_sycl_source_compilation_available(): + pytest.skip("SYCL source compilation extension not available") + if not q.get_sycl_device().can_compile("sycl"): + pytest.skip("SYCL source compilation not supported") + kb = dpc.create_kernel_bundle_from_sycl_source( + q, _local_accessor_sycl_source + ) + krn = kb.get_sycl_kernel("reverse_in_group") + + lws = 16 + gws = lws * 4 + x = np.arange(gws, dtype="i4") + x_usm = dpm.MemoryUSMDevice(x.nbytes, queue=q) + for k in (krn, make_kernel_of_unknown_origin(krn)): + q.memcpy(dest=x_usm, src=x, count=x.nbytes) + q.submit( + k, + [dpctl.LocalAccessor("i4", (lws,)), x_usm, ctypes.c_int(3)], + [gws], + [lws], + ).wait() + res = np.empty_like(x) + q.memcpy(dest=res, src=x_usm, count=x.nbytes) + expected = 3 * x.reshape(-1, lws)[:, ::-1].ravel() + assert np.array_equal(res, expected) + + # not checked by dpctl; the backend rejects the extra argument + with pytest.raises(dpctl.SyclKernelSubmitError): + q.submit_async( + k, + [ + dpctl.LocalAccessor("i4", (lws,)), + x_usm, + ctypes.c_int(3), + ctypes.c_int(3), + ], + [gws], + [lws], + ) + def test_submit_wrong_number_of_args(): try: @@ -367,3 +445,21 @@ def test_submit_wrong_number_of_args(): [x_usm, x_usm, x_usm, ctypes.c_int(2), ctypes.c_int(3)], [n], ) + + +def test_submit_wrong_number_of_args_opencl_source(): + try: + q = dpctl.SyclQueue("opencl") + except dpctl.SyclQueueCreationError: + pytest.skip("OpenCL queue could not be created") + kb = dpc.create_kernel_bundle_from_source( + q, "kernel void set(global int *a, int v) { a[0] = v; }" + ) + krn = kb.get_sycl_kernel("set") + assert krn.num_args == 2 + + x_usm = dpm.MemoryUSMDevice(4, queue=q) + with pytest.raises(ValueError): + q.submit_async(krn, [x_usm], [1]) + with pytest.raises(ValueError): + q.submit_async(krn, [x_usm, ctypes.c_int(1), ctypes.c_int(2)], [1]) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_interface.h index d8a69afb1d..fd0ba14dfe 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_interface.h @@ -45,7 +45,8 @@ DPCTL_C_EXTERN_C_BEGIN * @param KRef DPCTLSyclKernelRef pointer to a SYCL * interoperability kernel. * @return Returns the number of arguments for the interoperability - * kernel. + * kernel, or -1 if KRef is NULL or the number of arguments can + * not be queried. * @ingroup KernelInterface */ DPCTL_API diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index e035ea7000..f5bcfae8c5 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -53,8 +53,13 @@ size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef) } auto sycl_kernel = unwrap(KRef); - auto num_args = sycl_kernel->get_info(); - return static_cast(num_args); + try { + auto num_args = sycl_kernel->get_info(); + return static_cast(num_args); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return -1; + } } void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef) From 1dfebbbb01ccadf2aec2afa2c9d87685cac4e48e Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sun, 8 Nov 2026 14:41:07 -0800 Subject: [PATCH 2/3] hint at split local accessors in the argument count error A sycl::local_accessor of a kernel compiled from SYCL to SPIR-V can take several arguments; the error says so when a LocalAccessor is given, and the LocalAccessor docstring explains it. --- dpctl/_sycl_queue.pyx | 17 ++++++- .../local-accessor-free-function-kernel.spv | Bin 0 -> 2060 bytes dpctl/tests/test_sycl_kernel_submit.py | 43 ++++++++++++++++++ 3 files changed, 59 insertions(+), 1 deletion(-) create mode 100644 dpctl/tests/input_files/local-accessor-free-function-kernel.spv diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 52076a7230..a3ee3f66e0 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -175,6 +175,12 @@ cdef class LocalAccessor: Python class for specifying the dimensionality and type of a ``sycl::local_accessor``, to be used as a kernel argument type. + A ``LocalAccessor`` is passed as a single kernel argument. A + ``sycl::local_accessor`` parameter of a kernel compiled from SYCL to + SPIR-V can take several arguments, so such a kernel can not be given a + ``LocalAccessor``; create it with + :func:`dpctl.compiler.create_kernel_bundle_from_sycl_source` instead. + Args: dtype (str): the data type of the local memory. @@ -1574,10 +1580,19 @@ cdef class SyclQueue(_SyclQueue): if kernel._origin == _INTEROP_ORIGIN: nKA = kernel.num_args if len(args) != nKA: - raise ValueError( + msg = ( f"Kernel '{kernel.get_function_name()}' takes {nKA} " f"argument(s), got {len(args)}." ) + if len(args) < nKA and any( + [isinstance(arg, LocalAccessor) for arg in args] + ): + msg += ( + " A sycl::local_accessor parameter of a kernel " + "compiled from SYCL to SPIR-V can take several " + "arguments, see the dpctl.LocalAccessor documentation." + ) + raise ValueError(msg) # Allocate the arrays to be sent to DPCTLQueue_Submit kargs = malloc(len(args) * sizeof(void*)) diff --git a/dpctl/tests/input_files/local-accessor-free-function-kernel.spv b/dpctl/tests/input_files/local-accessor-free-function-kernel.spv new file mode 100644 index 0000000000000000000000000000000000000000..20bc522169e4c8774f25c37df1c65b926f02c80f GIT binary patch literal 2060 zcmZ{l+fEZv6oz*x=Q9H00dWvT5k>I;rbc5Vp%O_@N$_$~%Rm~&wwV?pUKuaE_7S}G z&ZqD>d<<_;|8FxZ)>1cFS?7PR!=93Do9S{c;|5))d*F`Gh--8065$;w?@W1L%8$mK z>vdgj^Tk&9{MoB8dRL8#^{`N@?Z1r_)2Gag_(nAf26tWQFKF4Y^_6=-j9% z3~PtQINXbpaukPgrC5l=La`XtYLz6l`Z|bLv94Q`)gLu}`IM25+dNLSlpKUl_e*ho zt-KDOZ#N*!BT7pJLsV42J0ol0Rqv;@p+~J3SN1k;})!1L^F6yU|DG=}GPE5FhnUN*WTk zTBj7vH-iniUmP90VR7CQHuMp3W&_T&wsuOKJHRK_h`2G`z*~N(Jnkls+DvB$jXL)< z>bwc`In@Nt1<{gFxR2jE4)-t1&;7hbWNP4@namq8nSCQuAG`gj8hP7~^z|~cb0X?w zHi77+G-SSUe4MX}#dc9dP59XFH^oxtvb;G|T8+|w9&&Y<4WB8XPgO3k< z!?+@uSv6zK%7%|)7*{2Ofe(z2X+PH_gOA_p<~IU+Afj&cRmC%_>mufbpE}t8ya+pW zuwP=?1K$g~U>*0=Me1V7_ v(7&fW-;+#Dd_yZ~t@kBUD>-Icl8l{r>vdT&cI$z@=pDP+j$B6n`!ntj?ZnQ~ literal 0 HcmV?d00001 diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index 81e700aaf2..2af43b2514 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -447,6 +447,49 @@ def test_submit_wrong_number_of_args(): ) +# The kernel in local-accessor-free-function-kernel.spv was generated with +# `icpx -fsycl -save-offload-code=` (DPC++ 2026.2) from: +# +# #include +# +# namespace syclexp = sycl::ext::oneapi::experimental; +# namespace syclext = sycl::ext::oneapi; +# +# extern "C" SYCL_EXTERNAL +# SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +# void local_accessor_kernel(sycl::local_accessor lmem, int *x) { +# auto item = syclext::this_work_item::get_nd_item<1>(); +# size_t lid = item.get_local_linear_id(); +# lmem[lid] = x[item.get_global_linear_id()]; +# sycl::group_barrier(item.get_group()); +# x[item.get_global_linear_id()] = +# lmem[item.get_local_range(0) - 1 - lid]; +# } +# +# int main() { return 0; } +def test_submit_local_accessor_arg_spirv_from_sycl(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Could not create default queue") + fn = get_spirv_abspath("local-accessor-free-function-kernel.spv") + with open(fn, "br") as f: + spirv_bytes = f.read() + kb = dpc.create_kernel_bundle_from_spirv(q, spirv_bytes) + krn = kb.get_sycl_kernel("__sycl_kernel_local_accessor_kernel") + # the local accessor takes several arguments in this SPIR-V + assert krn.num_args > 2 + + x_usm = dpm.MemoryUSMDevice(64, queue=q) + with pytest.raises(ValueError, match="dpctl.LocalAccessor documentation"): + q.submit_async( + krn, [dpctl.LocalAccessor("i4", (16,)), x_usm], [16], [16] + ) + with pytest.raises(ValueError) as e: + q.submit_async(krn, [x_usm], [16], [16]) + assert "LocalAccessor" not in str(e.value) + + def test_submit_wrong_number_of_args_opencl_source(): try: q = dpctl.SyclQueue("opencl") From 40dcdf28ea33533ee397a02003862d38ad7c1a89 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sun, 8 Nov 2026 15:33:13 -0800 Subject: [PATCH 3/3] document passing a split local accessor A LocalAccessor works for a kernel compiled from SYCL to SPIR-V when the accessor's ranges are passed after it; say so in the docstring and the error, and test it. --- dpctl/_sycl_queue.pyx | 21 +++++++++++++-------- dpctl/tests/test_sycl_kernel_submit.py | 24 +++++++++++++++++------- 2 files changed, 30 insertions(+), 15 deletions(-) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index a3ee3f66e0..3918e3262b 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -175,11 +175,15 @@ cdef class LocalAccessor: Python class for specifying the dimensionality and type of a ``sycl::local_accessor``, to be used as a kernel argument type. - A ``LocalAccessor`` is passed as a single kernel argument. A - ``sycl::local_accessor`` parameter of a kernel compiled from SYCL to - SPIR-V can take several arguments, so such a kernel can not be given a - ``LocalAccessor``; create it with - :func:`dpctl.compiler.create_kernel_bundle_from_sycl_source` instead. + A ``LocalAccessor`` is passed as a single kernel argument, the pointer to + the local memory. A kernel compiled from SYCL to SPIR-V may expect a + ``sycl::local_accessor`` parameter as several arguments instead: the + pointer, followed by the accessor's access range, memory range and + offset. Pass the ``LocalAccessor`` as the first of them and the others + explicitly, e.g. for a one-dimensional accessor of ``n`` elements:: + + [dpctl.LocalAccessor("i4", (n,)), ctypes.c_size_t(n), + ctypes.c_size_t(n), ctypes.c_size_t(0), ...] Args: dtype (str): @@ -1588,9 +1592,10 @@ cdef class SyclQueue(_SyclQueue): [isinstance(arg, LocalAccessor) for arg in args] ): msg += ( - " A sycl::local_accessor parameter of a kernel " - "compiled from SYCL to SPIR-V can take several " - "arguments, see the dpctl.LocalAccessor documentation." + " A kernel compiled from SYCL to SPIR-V may expect a " + "sycl::local_accessor parameter as several arguments, " + "of which a dpctl.LocalAccessor is only the first; " + "see the dpctl.LocalAccessor documentation." ) raise ValueError(msg) diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index 2af43b2514..ae0ac1a2ce 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -477,18 +477,28 @@ def test_submit_local_accessor_arg_spirv_from_sycl(): spirv_bytes = f.read() kb = dpc.create_kernel_bundle_from_spirv(q, spirv_bytes) krn = kb.get_sycl_kernel("__sycl_kernel_local_accessor_kernel") - # the local accessor takes several arguments in this SPIR-V - assert krn.num_args > 2 + # the local accessor is the pointer, access range, memory range and + # offset in this SPIR-V + assert krn.num_args == 5 - x_usm = dpm.MemoryUSMDevice(64, queue=q) + lws = 16 + gws = lws * 4 + x = np.arange(gws, dtype="i4") + x_usm = dpm.MemoryUSMDevice(x.nbytes, queue=q) + q.memcpy(dest=x_usm, src=x, count=x.nbytes) + la = dpctl.LocalAccessor("i4", (lws,)) with pytest.raises(ValueError, match="dpctl.LocalAccessor documentation"): - q.submit_async( - krn, [dpctl.LocalAccessor("i4", (16,)), x_usm], [16], [16] - ) + q.submit_async(krn, [la, x_usm], [gws], [lws]) with pytest.raises(ValueError) as e: - q.submit_async(krn, [x_usm], [16], [16]) + q.submit_async(krn, [x_usm], [gws], [lws]) assert "LocalAccessor" not in str(e.value) + sz = ctypes.c_size_t + q.submit(krn, [la, sz(lws), sz(lws), sz(0), x_usm], [gws], [lws]).wait() + res = np.empty_like(x) + q.memcpy(dest=res, src=x_usm, count=x.nbytes) + assert np.array_equal(res, x.reshape(-1, lws)[:, ::-1].ravel()) + def test_submit_wrong_number_of_args_opencl_source(): try: