[ndgrigorian]

This PR fixes the check on the number of kernel arguments added to SyclQueue.submit and SyclQueue.submit_async in gh-2382. It compares the number of arguments given with info::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 a sycl::local_accessor takes 4 arguments, so such a kernel can not be submitted with a dpctl.LocalAccessor:

ValueError: Kernel 'reverse_in_group' takes 6 argument(s), got 3.

dpctl.compiler.SyclKernel now 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 through SyclKernel_Make, are not checked. On OpenCL, the backend rejects a wrong number of arguments on submission, raising SyclKernelSubmitError, while on Level Zero too few arguments are not detected

A kernel compiled from SYCL to SPIR-V may also expect a sycl::local_accessor parameter as several arguments: the local memory pointer, followed by the accessor's access range, memory range and offset. A dpctl.LocalAccessor provides 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 a dpctl.LocalAccessor, and the dpctl.LocalAccessor docstring shows how to pass them:

ValueError: Kernel '__sycl_kernel_local_accessor_kernel' takes 5 argument(s), got 2. 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.

DPCTLKernel_GetNumArgs also no longer aborts the process when the query throws, returning -1, and SyclKernel.get_num_args raises RuntimeError in that case

• Have you provided a meaningful PR description?
• Have you added a test, reproducer or referred to an issue with a reproducer?
• Have you tested your changes locally for CPU and GPU devices?
• Have you made sure that new changes do not introduce compiler warnings?
• Have you checked performance impact of proposed changes?
• Have you added documentation for your changes, if necessary?
• Have you added your changes to the changelog?
• If this PR is a work in progress, are you opening the PR as a draft?

[github-actions]

View rendered docs @ https://intelpython.github.io/dpctl/pulls/2425/index.html

[coveralls]

coverage: 75.112% (-0.03%) from 75.137% — fix-kernel-arg-count-check into master