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

[gmlueck]

I haven't read P3312. Would this solve:

• Template argument deduction, or does it only solve the overload resolution case?
• The overload cases where Args are not exactly the same types as any of the overloads, but they are convertible to one of the overloads.

[gmlueck]

This part of the comment makes it sound like __builtin_sycl_launch_kernel only needs to be supported by the device compiler, but that's not true, is it? In order to use SYCL_EXT_ONEAPI_KERNEL_FUNCTION, doesn't the host compiler need to support that builtin?

[koparasy]

You are correct. It reads weird and it is incorrect. I will fix it.

[gmlueck]

Suggested change

	This macro requires support from the device compiler. An implementation that
	This macro requires support from the host compiler. An implementation that

I think ...

[koparasy]

You require support from both (host and device). Host side needs to parse this and not throw an error. But host side does not need to define the function, only instantiate the declaration. I will need to think a little more on how I should phrase this better.

[gmlueck]

I suggest making a new section for this example. That way, we can easily send a link to this section whenever someone asks about this feature via email, etc.

Also, rewrite the examples below to be full SYCL applications. It should be possible to cut / paste the example into a text file and then compile and run it.

[koparasy]

Will do.

[gmlueck]

I'm not sure this is necessary. We don't normally bump the feature-test macro for experimental extensions. @rolandschulz do you think your code will need the feature-test macro?

[koparasy]

No strong opinions here. I will do whatever the actual process requires.

[gmlueck]

This sentence could be improved:

• "args are used only to resolve the kernel" is not strictly correct. They are also passed to the launch.
• "are not evaluated more than once" -> "are evaluated exactly once"

Maybe:

The args are used to resolve the kernel name and are also passed when launching the kernel. Each argument is evaluated exactly once.

[gmlueck]

Rather than defining a separate macro, could applications just test whether SYCL_EXT_ONEAPI_KERNEL_FUNCTION is defined?

[koparasy]

I think they could, yes.

[tahonermann]

I haven't read P3312. Would this solve:

* Template argument deduction, or does it only solve the overload resolution case?

* The overload cases where `Args` are not exactly the same types as any of the overloads, but they are convertible to one of the overloads.

Yes, it would handle both of these. See P3312R1 section 5.1 for an as-if like description of how the overload set type would work.

At the moment, no further work on P3312 is planned within WG21; see cplusplus/papers#1963. Despite some good support for the paper, the number of people that voted neutral on the latest poll was seen as insufficient motivation for the paper to move forward. Additional motivation could revive the paper.