Check the number of kernel arguments only for interoperability kernels

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)

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
@@ -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), ...]
+
     Args:
         dtype (str):
             the data type of the local memory.
@@ -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.
@@ -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:
+            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*))
@@ -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.

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)
 

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 == <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.
@@ -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)

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",
 ]

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())