[SYCL][FFK] Add SYCL_EXT_ONEAPI_KERNEL_FUNCTION bare-name kernel launch

[koparasy]

Add an experimental way to launch a free function kernel by naming it directly,
letting the compiler deduce the kernel's template arguments / resolve the
overload from the launch arguments, instead of spelling a fully-resolved
kernel_function<Func>:

      nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(iota, 3.14f, ptr));
      single_task(q, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(store42, ptr));

The existing free function kernel launch API (kernel_function, nd_launch,
single_task) is unchanged; this is a thin, optional ergonomic layer on top.
Mechanism:

• __builtin_sycl_launch_kernel(name, args...): a front-end builtin
  (CustomTypeChecking, like the sibling __builtin_sycl_is_kernel family) that
  resolves the bare/overloaded/templated kernel name against the argument types
  using ordinary C++ overload resolution / template argument deduction, and
  evaluates to a pointer to the chosen specialization.
• SYCL_EXT_ONEAPI_KERNEL_FUNCTION macro (free_function_traits.hpp): fills the
  kernel_function selector of a launch with the deduced specialization and
  forwards the launch arguments. Codegen is identical to the explicit
  kernel_function form (the launch arguments are consumed only to drive
  deduction and are evaluated exactly once). It is gated on
  __has_builtin(__builtin_sycl_launch_kernel).

Non-deducible template parameters (a non-type parameter appearing in no function
parameter, a type parameter used only in a non-deduced context, ...) must be
spelled explicitly, e.g. SYCL_EXT_ONEAPI_KERNEL_FUNCTION((fill<1>), p, v). This
is standard C++ template argument deduction behavior.

Documented in the sycl_ext_oneapi_free_function_kernels extension (new
"Launching a kernel by name" section, feature-test value 2, examples). The spec
describes the macro's effect and leaves its return type unspecified so other
implementations may lower it differently.

Dependent-context note: using the macro inside a template with dependent
arguments would classify a deferred builtin call whose callee still has the
BuiltinFn placeholder type, tripping a front-end assertion in
CallExpr::getCallReturnType. This is a general Clang defect (reproducible with
other custom-type-checked builtins, fixed upstream by llvm/llvm-project's
[Clang] Fix assertion failure when classifying a dependent call to a builtin.
Until that fix reaches this tree, the macro wraps the builtin in a unary '+' so
the argument is a UnaryOperator rather than the raw CallExpr, which is standard
C++ (identity on a function pointer) and avoids the classification path. The
sentinel test flips when the underlying fix lands, signalling the workaround can
be removed.

Co-Authored-By: Claude Opus 4.8 (1M context) noreply@anthropic.com

[rolandschulz]

Very nice! Only reviewed spec and tests from a user requirements perspective.

[tahonermann]

I haven't reviewed this closely yet, but added a couple of superficial comments.

I'm generally opposed to what this is trying to do. I understand the desire to free the user from having to write casts, but they can choose to avoid that by using unique names for their kernels (in most cases, with some exceptions for generic code where functionality like this is sorely missing from the core C++ language). I would prefer we provide a general feature for this rather than something specific to this very narrow case of resolving a SYCL free function kernel from an overload set.

I mentioned this in one of my other comments, but this functionality should not be compared to CUDA's kernel call expressions. The macro is inelegant and suffers from the usual macro problems. The comparison to CUDA is a pale one at best.

[tahonermann]

I don't think this is a good name for what this builtin does. It doesn't actually launch a kernel; all it does is select an overload from an overload set based on a set of arguments. Other name suggestions:

• __builtin_sycl_declcall (name stolen from https://www.open-std.org/jtc1/sc22/wg21/docs/papers/2024/p2825r1.html).
• __builtin_sycl_kernel_selector
• __builtin_sycl_overload_resolver

[tahonermann]

Likewise with respect to my comment about the name of the builtin function, this name doesn't reflect what the function really does. It is misleadingly similar to other functions in this same file that actually are involved in launching a kernel (e.g., BuildSYCLKernelLaunchCallArgs()).

I find the comparison with CUDA misleading as well. The CUDA kernel call expression does actually launch a kernel (by implicitly calling cudaConfigureCall() or whatever the current name of that function is).

[koparasy]

I haven't reviewed this closely yet, but added a couple of superficial comments.

I'm generally opposed to what this is trying to do. I understand the desire to free the user from having to write casts, but they can choose to avoid that by using unique names for their kernels (in most cases, with some exceptions for generic code where functionality like this is sorely missing from the core C++ language). I would prefer we provide a general feature for this rather than something specific to this very narrow case of resolving a SYCL free function kernel from an overload set.

I mentioned this in one of my other comments, but this functionality should not be compared to CUDA's kernel call expressions. The macro is inelegant and suffers from the usual macro problems. The comparison to CUDA is a pale one at best.

I'm ok bikeshedding the builtin name and rewording the docs to drop the CUDA comparison.

The declcall direction may well be valid,. I'm not yet convinced it removes the need for the macro, though. As I understand it, declcall returns a pointer to the resolved function. Like the current builtin, it still needs the call arguments to drive overload resolution / deduction (it presumably parses the call expression itself, in an unevaluated context). But the result is just a function pointer: to actually launch, the arguments still have to be passed to the kernel. So, the launch site needs the arguments in two roles, once to resolve the kernel, once to pass to it, and writing them a single time is what requires the macro.

Concretely, nd_launch(q, ndr, kernel_function<declcall(k(args))>, args) works but repeats args; folding that into one spelling is exactly what the macro does. I believe that the arguments can't be deferred into the launch function and resolved there: a bare overloaded/templated kernel name has and can't be bound to a parameter, so resolution must happen at the call site where the name and its arguments meet.

The current approach is a table builtin with custom type checking, no custom parsing, and the Sema side is a single handler that reuses the normal call-building machinery (BuildCallExpr + read the resolved callee). A declcall-style operator needs a dedicated parse path for the unevaluated call and correspondingly more Sema. Resolution still happens in Sema either way, but as a first-class operator it likely wants its own AST node with the usual plumbing (template instantiation, constant evaluation, serialization). Overload resolution itself is the same work in both; declcall adds front-end surface around it rather than removing any. And the genuinely subtle part. Classifying a not-yet-resolved call while deducing the enclosing template exists in both, so declcall doesn't make clang simpler.

To put it plainly: declcall addresses the resolution primitive. The part that's already solved here. The hard part isn't the resolver; it's designing a launch API that passes the arguments once without a macro, and I don't yet see that without either a macro or a new language-level launch construct. I'd genuinely like to be convinced otherwise., if you can sketch a launch API that avoids the macro (without requiring a new language-level launch expression) and avoids us redesigning all the APIs and implementations.

[tahonermann]

The declcall direction may well be valid,. I'm not yet convinced it removes the need for the macro, though.

I agree that declcall doesn't have a solution for avoiding repetition of the arguments (but at least they would be unevaluated within the declcall operator). A solution that avoids that repetition would be the ability to pass an overload set as a template argument along the lines of what has been proposed in these papers:

• P3312: Overload Set Types
• P1170: Overload sets as function parameters
• P0119: Overload sets as function arguments

See section 3.3 of P3312 for a discussion of the limitations of declcall that you observed.

I'm not aware of any implementation experience for those proposals and I don't have a good sense of how difficult an implementation would be.

Classifying a not-yet-resolved call while deducing the enclosing template exists in both, so declcall doesn't make clang simpler.

I agree. I'm concerned about the user experience. I'm not particularly concerned about implementation other than that I don't want to have to maintain something that is likely to be obsoleted by better language features. I would rather users use unique names and write static casts until that better language feature is available (with the acknowledgement that those workarounds are pretty limited in generic programming contexts).

I'd genuinely like to be convinced otherwise., if you can sketch a launch API that avoids the macro (without requiring a new language-level launch expression) and avoids us redesigning all the APIs and implementations.

P3312 would allow nd_launch() to be declared something like the following (where I've injected my preference for the overload set to be passed as an explicit template argument, but the sycl::kernel_function<...> approach would work too).

template<auto KernelOverloadSet, int Dimensions, typename... Args>
void nd_launch(sycl::queue &q, sycl::nd_range<Dimensions> ndr, Args&&... args);

Kernel launch would then look like:

SYCL_KERNEL void my_kernel(int ip*, int n) { ... }
SYCL_KERNEL void my_kernel(float fp*, int n) { ... }

int n = 42;
int *ip = sycl::malloc_shared<int>(n, q);
sycl::nd_launch<my_kernel>(q, ndr, ip, n);

Or with the sycl::kernel_function variable template approach:

sycl::nd_launch(q, ndr, sycl::kernel_function<my_kernel>, ip, n);

That still isn't as nice as the CUDA kernel launch expression, but I think it would still be a significant improvement over what has been proposed so far.

An aspect that hasn't been discussed (to my knowledge) is how any of these approaches work in the context of sycll:get_kernel_id<>(). The declcall approach works naturally there and declcall or something very much like it would still be needed for the overload set approach to actually select an overload. For example, with the above examples, the user is still going to need to be able to do something like this:

sycl::kernel_id kid = sycl::get_kernel_id<declcall(my_kernel, std::declval<int*>(), 0)>();