Skip to content

Check the number of kernel arguments only for interoperability kernels #2425

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
ndgrigorian wants to merge 3 commits into
base: master
Choose a base branch
from
Open

Check the number of kernel arguments only for interoperability kernels #2425

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
b7f435b
check kernel argument count only for interop kernels
ndgrigorian Nov 8, 2026
1dfebbb
hint at split local accessors in the argument count error
ndgrigorian Nov 8, 2026
40dcdf2
document passing a split local accessor
ndgrigorian Nov 8, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand All @@ -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)
Expand Down
54 changes: 43 additions & 11 deletions dpctl/_sycl_queue.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -174,6 +175,16 @@ 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, 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), ...]
Comment on lines +178 to +186

@vlad-perevezentsev vlad-perevezentsev Nov 10, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For an N-D accessor, each range and the offset is one argument of N size_t, not N arguments. A SPIR-V kernel with sycl::local_accessor<int, 2> and int * has 5 arguments, and it gives the right result on GPU and CPU with:

r = lambda a, b: dpctl.RawKernelArg(np.array([a, b], dtype=np.uintp))
args = [dpctl.LocalAccessor("i4", (4, 4)), r(4, 4), r(4, 4), r(0, 0), x_usm]

Could the docstring say this?


Args:
dtype (str):
the data type of the local memory.
Expand Down Expand Up @@ -1540,8 +1551,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.
Expand All @@ -1564,13 +1580,24 @@ 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 <size_t>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:

@vlad-perevezentsev vlad-perevezentsev Nov 10, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

SYCL source kernels are no longer checked at all. On Level Zero a missing argument now runs the kernel silently, while master raises ValueError, and for a kernel without a local accessor the count is right:

import numpy as np, dpctl, dpctl.compiler as dpc, dpctl.memory as dpm

src = """
#include <sycl/sycl.hpp>
namespace syclexp = sycl::ext::oneapi::experimental;
extern "C" SYCL_EXTERNAL
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
void fill(int *x, int v) {
    auto it = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
    x[it.get_global_linear_id()] = v;
}
"""
q = dpctl.SyclQueue()
krn = dpc.create_kernel_bundle_from_sycl_source(q, src).get_sycl_kernel("fill")
x = dpm.MemoryUSMDevice(64 * 4, queue=q)
q.submit(krn, [x], [64], [16]).wait()  # `v` is missing
res = np.empty(64, dtype="i4")
q.memcpy(res, x, res.nbytes)
print(res[:4])  # Level Zero GPU: [0 0 0 0]; OpenCL CPU: SyclKernelSubmitError; master: ValueError

Could the check stay for SYCL source kernels when no LocalAccessor is passed, or is the count unreliable there too?

nKA = kernel.num_args
if <size_t>len(args) != nKA:
msg = (
f"Kernel '{kernel.get_function_name()}' takes {nKA} "
f"argument(s), got {len(args)}."
)
if <size_t>len(args) < nKA and any(
[isinstance(arg, LocalAccessor) for arg in args]
):
msg += (
" 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)

# Allocate the arrays to be sent to DPCTLQueue_Submit
kargs = <void**>malloc(len(args) * sizeof(void*))
Expand Down Expand Up @@ -1718,8 +1745,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.
Expand Down
17 changes: 13 additions & 4 deletions dpctl/compiler/_compiler.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

@vlad-perevezentsev vlad-perevezentsev Nov 10, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The new _origin field makes PySyclKernelObject 8 bytes larger. Extensions built against an earlier dpctl that cimport dpctl or include dpctl4pybind11.hpp (all of dpnp) now get a RuntimeWarning on import, and fail with warnings as errors:

# dpnp built against dpctl master, run with this branch
import warnings
warnings.simplefilter("error", RuntimeWarning)
import dpnp  # RuntimeWarning: dpctl.compiler._compiler.SyclKernel size changed ... Expected 40 from C header, got 48

Could interop kernels be a subclass of SyclKernel without new fields instead, so the struct and the _create vtable slot stay the same?

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 [
Expand All @@ -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)

Expand Down
39 changes: 27 additions & 12 deletions dpctl/compiler/_compiler.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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):
Expand All @@ -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 == <size_t>-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.
Expand Down Expand Up @@ -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):
Expand Down Expand Up @@ -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
)
Expand All @@ -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
)
Expand Down Expand Up @@ -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:
Expand Down Expand Up @@ -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(
Expand Down Expand Up @@ -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,
Expand Down Expand Up @@ -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)
2 changes: 2 additions & 0 deletions dpctl/tests/helper/__init__.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@
has_cpu,
has_gpu,
has_sycl_platforms,
make_kernel_of_unknown_origin,
)

__all__ = [
Expand All @@ -30,4 +31,5 @@
"has_gpu",
"has_sycl_platforms",
"get_queue_or_skip",
"make_kernel_of_unknown_origin",
]
21 changes: 21 additions & 0 deletions dpctl/tests/helper/_helper.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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())
Binary file added dpctl/tests/input_files/local-accessor-free-function-kernel.spv
View file
Open in desktop
Binary file not shown.
Loading
Loading