Repository navigation
Check the number of kernel arguments only for interoperability kernels - #2425
ndgrigorian wants to merge 3 commits into
Conversation
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.
c562604 to
b7f435b
Compare
|
View rendered docs @ https://intelpython.github.io/dpctl/pulls/2425/index.html |
|
coverage: 75.112% (-0.03%) from 75.137% — fix-kernel-arg-count-check into master |
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.
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.
| ''' | ||
| cdef DPCTLSyclKernelRef _kernel_ref | ||
| cdef str _function_name | ||
| cdef int _origin |
There was a problem hiding this comment.
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 48Could interop kernels be a subclass of SyclKernel without new fields instead, so the struct and the _create vtable slot stay the same?
| f"argument(s), got {len(args)}." | ||
| ) | ||
| # provided, so check it here for the kernels it defines that number for | ||
| if kernel._origin == _INTEROP_ORIGIN: |
There was a problem hiding this comment.
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: ValueErrorCould the check stay for SYCL source kernels when no LocalAccessor is passed, or is the count unreliable there too?
| 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), ...] |
There was a problem hiding this comment.
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?
This PR fixes the check on the number of kernel arguments added to
SyclQueue.submitandSyclQueue.submit_asyncin gh-2382. It compares the number of arguments given withinfo::kernel::num_args, which SYCL defines only for kernels in kernel bundles created through backend interoperability. For a kernel compiled from SYCL source, DPC++ reports the backend's count instead, in which asycl::local_accessortakes 4 arguments, so such a kernel can not be submitted with adpctl.LocalAccessor:dpctl.compiler.SyclKernelnow records the kind of kernel bundle it comes from, and the number of arguments is only checked for kernels from kernel bundles created from OpenCL C source or SPIR-V. The arguments of other kernels, including those passed to dpctl throughSyclKernel_Make, are not checked. On OpenCL, the backend rejects a wrong number of arguments on submission, raisingSyclKernelSubmitError, while on Level Zero too few arguments are not detectedA kernel compiled from SYCL to SPIR-V may also expect a
sycl::local_accessorparameter as several arguments: the local memory pointer, followed by the accessor's access range, memory range and offset. Adpctl.LocalAccessorprovides only the pointer, so the others have to be passed explicitly after it. For kernels from a SPIR-V bundle, the error now says so when the arguments include adpctl.LocalAccessor, and thedpctl.LocalAccessordocstring shows how to pass them:DPCTLKernel_GetNumArgsalso no longer aborts the process when the query throws, returning -1, andSyclKernel.get_num_argsraisesRuntimeErrorin that case