diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td index bffdb6b30b46e..541f3abecaf0b 100644 --- a/clang/include/clang/Basic/Builtins.td +++ b/clang/include/clang/Basic/Builtins.td @@ -5317,6 +5317,21 @@ def SYCLIsNDRangeKernel : LangBuiltin<"SYCL_LANG"> { let Prototype = "bool(...)"; } +// Resolves a (possibly overloaded / templated) free-function-kernel name +// against the provided argument types using ordinary C++ overload resolution +// and template-argument deduction, and evaluates to a pointer to the resolved +// specialization. Used by the SYCL_KERNEL(name, args...) launch macro to fill +// the `kernel_function` non-type template parameter of the enqueue +// functions from a bare kernel name, e.g. +// nd_launch(q, range, SYCL_KERNEL(iota, 3.14f, p)); +// Arg 0 is the unresolved kernel name; the remaining args are the launch +// arguments (consumed only to drive deduction, never evaluated). +def SYCLLaunchKernel : LangBuiltin<"SYCL_LANG"> { + let Spellings = ["__builtin_sycl_launch_kernel"]; + let Attributes = [NoThrow, CustomTypeChecking]; + let Prototype = "void(...)"; +} + // HLSL def HLSLAddUint64: LangBuiltin<"HLSL_LANG"> { let Spellings = ["__builtin_hlsl_adduint64"]; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 568a18e2c375c..4d1c132a8ccfb 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13955,6 +13955,12 @@ def err_sycl_add_ir_attribute_invalid_value : Error< def err_sycl_add_ir_attribute_invalid_filter : Error< "initializer list in the first argument of %0 must contain only string " "literals">; +// Bare-name free-function-kernel launch (__builtin_sycl_launch_kernel, used by +// the SYCL_EXT_ONEAPI_KERNEL_FUNCTION macro). +def err_sycl_launch_kernel_not_free_function : Error< + "the function resolved by '__builtin_sycl_launch_kernel' is not a SYCL free " + "function kernel; it must be declared with 'sycl-nd-range-kernel' or " + "'sycl-single-task-kernel' compile time properties">; def warn_sycl_old_and_new_kernel_attributes : Warning< "kernel has both attribute %0 and kernel properties; if the kernel " "properties contains the property \"%1\" it will be ignored">, diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index c05886f988a93..39a129c9e4ab4 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -391,6 +391,16 @@ class SemaSYCL : public SemaBase { void ProcessFreeFunction(FunctionDecl *FD); void finalizeFreeFunctionKernels(); + /// Handle the '__builtin_sycl_launch_kernel' builtin. Arg 0 of \p TheCall is + /// a (possibly overloaded / templated) SYCL free-function-kernel name; the + /// remaining args are the launch arguments. Resolves the name against those + /// argument types using ordinary C++ overload resolution / template argument + /// deduction, verifies the result is a free function kernel, and evaluates to + /// a pointer to the resolved specialization (usable as the `auto *Func` + /// non-type template parameter of the enqueue-function launch path). Used by + /// the SYCL_KERNEL(name, args...) launch macro. + ExprResult BuildSYCLLaunchKernelCall(CallExpr *TheCall); + /// Get the number of fields or captures within the parsed type. ExprResult ActOnSYCLBuiltinNumFieldsExpr(ParsedType PT); ExprResult BuildSYCLBuiltinNumFieldsExpr(SourceLocation Loc, diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index 11fe2a3cb774e..0319a899b0a12 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -4190,6 +4190,9 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID, break; } + + case Builtin::BI__builtin_sycl_launch_kernel: + return SYCL().BuildSYCLLaunchKernelCall(TheCall); } if (getLangOpts().HLSL && HLSL().CheckBuiltinFunctionCall(BuiltinID, TheCall)) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 62a610342fec0..b0ba24c9c1814 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -5917,6 +5917,106 @@ void SemaSYCL::processFreeFunctionDeclaration(const FunctionDecl *FD) { FreeFunctionDeclarations.insert(FD->getCanonicalDecl()); } +// Handle __builtin_sycl_launch_kernel(name, launch-args...). The +// SYCL_EXT_ONEAPI_KERNEL_FUNCTION launch macro expands to +// +// kernel_function<__builtin_sycl_launch_kernel(name, args...)>, args... +// +// so this builtin sits in the `kernel_function` non-type template +// parameter slot of an enqueue function. Its job is purely to resolve `name` +// (a bare, possibly overloaded / templated free-function-kernel name) to a +// concrete specialization using the launch argument types, and to evaluate to +// a pointer to it — a constant expression usable as the `auto *Func` NTTP. The +// same launch arguments are passed positionally to the enqueue function to +// perform the actual launch; the builtin never evaluates them. Mirrors the +// CUDA `<<<>>>` front-end approach (build a synthetic call, read the resolved +// callee), and reuses the real overload-resolution machinery so a bad launch +// produces ordinary diagnostics at the call site. +ExprResult SemaSYCL::BuildSYCLLaunchKernelCall(CallExpr *TheCall) { + if (TheCall->getNumArgs() < 1) { + // The builtin needs at least the kernel-name argument. This is unreachable + // through the SYCL_EXT_ONEAPI_KERNEL_FUNCTION macro (NAME is mandatory); + // the guard only protects the getArg(0) below for a direct zero-argument + // call. Reuse the generic builtin arg-count diagnostic (as the sibling + // __builtin_sycl_is_kernel family does) rather than a bespoke one. + Diag(TheCall->getBeginLoc(), diag::err_builtin_invalid_argument_count) << 1; + return ExprError(); + } + + Expr *KernelNameExpr = TheCall->getArg(0); + SmallVector LaunchArgs(TheCall->arguments().begin() + 1, + TheCall->arguments().end()); + + // Depend on template parameters: defer until instantiation. When any launch + // argument (or the kernel name) is dependent, Sema::BuildCallExpr has already + // created this builtin call with a dependent result type and skipped custom + // type checking, so we simply return it; this handler re-runs on the + // instantiated, concrete-typed call. The deferred call keeps its BuiltinFn + // placeholder callee; the SYCL_EXT_ONEAPI_KERNEL_FUNCTION macro wraps the + // builtin in a unary plus so the enclosing kernel_function<...> NTTP argument + // is a UnaryOperator (a plain prvalue that does not recurse into this + // dependent CallExpr) rather than the CallExpr itself, which is what keeps + // classification off CallExpr::getCallReturnType and avoids a front-end + // crash. + if (KernelNameExpr->isTypeDependent() || KernelNameExpr->isValueDependent() || + Expr::hasAnyTypeDependentArguments(LaunchArgs)) + return TheCall; + + // Build the synthetic call and let overload resolution + deduction run. This + // reuses the real C++ machinery so diagnostics (no viable overload, + // ambiguity, non-deducible template parameter) are emitted at the launch + // site for free. + ExprResult Call = SemaRef.BuildCallExpr(/*Scope=*/nullptr, KernelNameExpr, + KernelNameExpr->getBeginLoc(), + LaunchArgs, TheCall->getRParenLoc()); + if (Call.isInvalid()) + return ExprError(); + + auto *ResolvedCall = dyn_cast(Call.get()->IgnoreImpCasts()); + FunctionDecl *ResolvedFn = + ResolvedCall ? ResolvedCall->getDirectCallee() : nullptr; + if (!ResolvedFn) + // Overload resolution / deduction failed and already emitted diagnostics + // (BuildCallExpr can return a recovery expression rather than an invalid + // one), or the name did not resolve to a direct callee. Nothing to add. + return ExprError(); + + // Free-function-kernel verification is only meaningful during device + // compilation. On the host pass the SYCL driver includes the integration + // footer, whose attribute-less redeclaration of the kernel would make + // isFreeFunction() spuriously fail; the host pass only needs the resolved + // function pointer for the runtime launch. The device pass performs the real + // validation. + if (getLangOpts().SYCLIsDevice) { + if (!isFreeFunction(ResolvedFn)) { + Diag(KernelNameExpr->getBeginLoc(), + diag::err_sycl_launch_kernel_not_free_function); + Diag(ResolvedFn->getLocation(), diag::note_previous_declaration); + return ExprError(); + } + + // Force instantiation of the specialization so its definition is emitted + // and routed into the free-function-kernel collection/emission path exactly + // as the explicit kernel_function form does today. The synthetic Call + // above is discarded; no host call to the kernel is emitted. + SemaRef.MarkFunctionReferenced(KernelNameExpr->getBeginLoc(), ResolvedFn); + } + + // Evaluate to &ResolvedFn: an ordinary function-pointer constant expression + // (the builtin CallExpr itself is not constant-foldable), usable directly as + // the `auto *Func` non-type template parameter of kernel_function. The + // emitted SPIR-V kernel is therefore the real user function (no wrapper) and + // the runtime launch path is unchanged. + QualType FnPtrTy = getASTContext().getPointerType(ResolvedFn->getType()); + ExprResult FnRef = + SemaRef.BuildDeclRefExpr(ResolvedFn, ResolvedFn->getType(), VK_LValue, + KernelNameExpr->getBeginLoc()); + if (FnRef.isInvalid()) + return ExprError(); + return SemaRef.ImpCastExprToType(FnRef.get(), FnPtrTy, + CK_FunctionToPointerDecay); +} + void SemaSYCL::ProcessFreeFunction(FunctionDecl *FD) { if (isFreeFunction(FD)) { if (CheckFreeFunctionDiagnostics(SemaRef, FD)) diff --git a/clang/test/SemaSYCL/builtin_sycl_launch_kernel.cpp b/clang/test/SemaSYCL/builtin_sycl_launch_kernel.cpp new file mode 100644 index 0000000000000..dcc77f51d87e3 --- /dev/null +++ b/clang/test/SemaSYCL/builtin_sycl_launch_kernel.cpp @@ -0,0 +1,65 @@ +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device %s -verify + +// Tests __builtin_sycl_launch_kernel: a bare (overloaded / templated) SYCL free +// function kernel name is resolved against the launch argument types via +// ordinary C++ overload resolution / template argument deduction, and the +// builtin evaluates to a pointer to the resolved specialization — usable as the +// `auto *Func` non-type template parameter of the enqueue-function launch path. +// This is the front-end seam the SYCL_EXT_ONEAPI_KERNEL_FUNCTION launch macro is built on. + +#include "sycl.hpp" + +template +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void axpy(T *y, const T *x, T a, int n) { + for (int i = 0; i < n; ++i) + y[i] = a * x[i] + y[i]; +} + +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void ovl(int *p) {} // expected-note {{candidate function not viable}} + +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void ovl(float *p) {} // expected-note {{candidate function not viable}} + +// A plain function that is NOT a free-function kernel. +__attribute__((sycl_device)) +void not_a_kernel(int *p) {} // expected-note {{previous declaration is here}} + +// The builtin evaluates to a pointer to the resolved specialization, usable as +// an `auto *Func` non-type template parameter. +template struct kernel_function_s {}; + +void test() { + float *y = nullptr; + const float *x = nullptr; + int *ip = nullptr; + float *fp = nullptr; + + // Deduce T = float from the pointer arguments; result folds to '&axpy' + // and is usable as a non-type template argument. + kernel_function_s<__builtin_sycl_launch_kernel(axpy, y, x, 1.0f, 8)> k1; + (void)k1; + + // Pick the right overload by argument type. + kernel_function_s<__builtin_sycl_launch_kernel(ovl, ip)> k2; + (void)k2; + kernel_function_s<__builtin_sycl_launch_kernel(ovl, fp)> k3; + (void)k3; + + // No viable overload: no ovl(double*). + double *dp = nullptr; + // expected-error@+1 {{no matching function for call to 'ovl'}} + (void)__builtin_sycl_launch_kernel(ovl, dp); + + // Resolves to a real function that is not a free-function kernel. + // expected-error@+1 {{is not a SYCL free function kernel}} + (void)__builtin_sycl_launch_kernel(not_a_kernel, ip); + + // Missing kernel-name argument. + // expected-error@+1 {{builtin takes one argument}} + (void)__builtin_sycl_launch_kernel(); +} diff --git a/clang/test/SemaSYCL/builtin_sycl_launch_kernel_adversarial.cpp b/clang/test/SemaSYCL/builtin_sycl_launch_kernel_adversarial.cpp new file mode 100644 index 0000000000000..8f82841fbaaf6 --- /dev/null +++ b/clang/test/SemaSYCL/builtin_sycl_launch_kernel_adversarial.cpp @@ -0,0 +1,90 @@ +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device %s -verify + +// Adversarial coverage for __builtin_sycl_launch_kernel name resolution: the +// kernel-name operand must survive the full range of C++ ways a function name +// can be spelled, both non-dependent and (via the '+' form the SYCL_EXT_ONEAPI_KERNEL_FUNCTION +// macro emits) inside templates with dependent arguments. +// +// These lock in the behaviors that were only checked manually before: +// - qualified names (ns::ft) +// - explicit template-ids (ns::ft) +// - overloaded names resolved by argument type, including in a dependent +// context +// - mixed template parameters: a leading non-deducible NTTP given explicitly +// (Dim) plus a trailing deducible type parameter (T) +// - zero launch arguments + +#include "sycl.hpp" + +// expected-no-diagnostics + +template struct kernel_function_s {}; +template constexpr kernel_function_s kernel_function{}; + +namespace ns { +template +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void ft(T *p, int n) {} +} // namespace ns + +// Overload set, distinguished by argument type. +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void ovl(int *p) {} +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void ovl(float *p) {} + +// Mixed: Dim is a non-deducible NTTP (appears in no parameter), T is deducible. +template +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void mix(T *p, int n) {} + +// Zero-argument kernel. +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", "")]] +void noop() {} + +// Non-dependent uses. +void non_dependent() { + float *p = nullptr; + int *ip = nullptr; + + // Qualified name. + (void)kernel_function<+__builtin_sycl_launch_kernel(ns::ft, p, 8)>; + // Qualified explicit template-id. + (void)kernel_function<+__builtin_sycl_launch_kernel(ns::ft, p, 8)>; + // Overload resolved by argument type. + (void)kernel_function<+__builtin_sycl_launch_kernel(ovl, ip)>; + (void)kernel_function<+__builtin_sycl_launch_kernel(ovl, p)>; + // Mixed: Dim=2 explicit (leading, non-deducible), T deduced from p. + (void)kernel_function<+__builtin_sycl_launch_kernel(mix<2>, p, 8)>; + // Zero launch arguments. + (void)kernel_function<+__builtin_sycl_launch_kernel(noop)>; +} + +// Dependent uses (the '+' form the macro emits): resolution deferred to +// instantiation. These previously risked the front-end classification crash. +template +auto dep_qualified(T *p) { + return kernel_function<+__builtin_sycl_launch_kernel(ns::ft, p, 8)>; +} +template +auto dep_overload(T *p) { + return kernel_function<+__builtin_sycl_launch_kernel(ovl, p)>; +} +template +auto dep_mixed(T *p) { + return kernel_function<+__builtin_sycl_launch_kernel(mix, p, 8)>; +} + +void dependent() { + int *ip = nullptr; + float *fp = nullptr; + (void)dep_qualified(fp); // ns::ft + (void)dep_overload(ip); // ovl(int*) + (void)dep_overload(fp); // ovl(float*) + (void)dep_mixed<3>(fp); // mix<3, float> +} diff --git a/clang/test/SemaSYCL/builtin_sycl_launch_kernel_dependent.cpp b/clang/test/SemaSYCL/builtin_sycl_launch_kernel_dependent.cpp new file mode 100644 index 0000000000000..d036ca37ff8f7 --- /dev/null +++ b/clang/test/SemaSYCL/builtin_sycl_launch_kernel_dependent.cpp @@ -0,0 +1,40 @@ +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device %s -verify + +// Regression test: __builtin_sycl_launch_kernel used inside a template, with +// dependent launch arguments, must be deferred to instantiation rather than +// resolved eagerly. Using the raw builtin call as the `auto *Func` non-type +// template argument of kernel_function<...> would crash the front end while +// dependent: deducing the auto* parameter runs Expr::Classify on the builtin +// CallExpr, whose callee has the BuiltinFn placeholder type, and +// CallExpr::getCallReturnType asserts trying to castAs. +// +// The SYCL_EXT_ONEAPI_KERNEL_FUNCTION macro wraps the builtin in a unary plus, so the non-type +// template argument is a UnaryOperator (whose classification is a plain +// prvalue and does not recurse into the dependent CallExpr) instead of the +// CallExpr itself. That is what makes the dependent case well-formed with no +// front-end change; this test spells the same `+` form the macro emits. +// (Unary plus on a function pointer is the identity, so the deduced Func value +// is unchanged.) + +#include "sycl.hpp" + +template struct kernel_function_s {}; +template constexpr kernel_function_s kernel_function{}; + +template +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void axpy(T *y, const T *x, T a, int n) {} + +// Dependent-argument use inside a function template: deferred to instantiation. +template +auto select(T *y, const T *x, T a, int n) { + return kernel_function<+__builtin_sycl_launch_kernel(axpy, y, x, a, n)>; +} + +void ok() { + // expected-no-diagnostics + float *y = nullptr; + const float *x = nullptr; + (void)select(y, x, 1.0f, 8); // instantiates select, resolves axpy +} diff --git a/clang/test/SemaSYCL/builtin_sycl_launch_kernel_raw_dependent_crash.cpp b/clang/test/SemaSYCL/builtin_sycl_launch_kernel_raw_dependent_crash.cpp new file mode 100644 index 0000000000000..9dac3b47b90d7 --- /dev/null +++ b/clang/test/SemaSYCL/builtin_sycl_launch_kernel_raw_dependent_crash.cpp @@ -0,0 +1,41 @@ +// Boundary / sentinel test for the dependent-context crash. +// +// The SYCL_EXT_ONEAPI_KERNEL_FUNCTION macro wraps the builtin in a unary '+' so that, inside a +// template with dependent arguments, the enclosing kernel_function<...> auto* +// non-type template argument is a UnaryOperator rather than the raw builtin +// CallExpr. Classifying a UnaryOperator does not recurse into +// CallExpr::getCallReturnType, which is what avoids a front-end assertion +// (castAs() on the builtin's BuiltinFn placeholder callee). +// +// The '+' is a WORKAROUND for a latent front-end defect, not a fix: the RAW +// form (no '+') still crashes. This test pins that boundary. `not --crash` +// asserts clang aborts on the raw form. If a future clang change fixes the +// underlying classification / getCallReturnType path, clang will NO LONGER +// crash here, `not --crash` will fail, and this test will start failing -- +// which is the signal to revisit whether the '+' workaround is still needed +// (and whether the underlying fix should be upstreamed). +// +// RUN: not --crash %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device \ +// RUN: -fsyntax-only %s + +#include "sycl.hpp" + +template struct kernel_function_s {}; +template constexpr kernel_function_s kernel_function{}; + +template +__attribute__((sycl_device)) +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 1)]] +void axpy(T *y, const T *x, T a, int n) {} + +template +auto select(T *y, const T *x, T a, int n) { + // RAW builtin (no unary '+') as the auto* NTTP, in a dependent context. + return kernel_function<__builtin_sycl_launch_kernel(axpy, y, x, a, n)>; +} + +void use() { + float *y = nullptr; + const float *x = nullptr; + (void)select(y, x, 1.0f, 8); +} diff --git a/clang/test/SemaSYCL/kernel_function_macro_supported.cpp b/clang/test/SemaSYCL/kernel_function_macro_supported.cpp new file mode 100644 index 0000000000000..5679e91b7f1dc --- /dev/null +++ b/clang/test/SemaSYCL/kernel_function_macro_supported.cpp @@ -0,0 +1,26 @@ +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsyntax-only %s + +// Under the SYCL device compiler the __builtin_sycl_launch_kernel builtin is +// available, so the SYCL_EXT_ONEAPI_KERNEL_FUNCTION macro is enabled and its +// capability macro SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED is defined. This +// mirrors the __has_builtin gate the extension header uses so applications can +// detect the feature portably. + +#ifndef __has_builtin +#define __has_builtin(x) 0 +#endif + +#if !__has_builtin(__builtin_sycl_launch_kernel) +#error "__builtin_sycl_launch_kernel should be available under -fsycl" +#endif + +#if __has_builtin(__builtin_sycl_launch_kernel) +#define SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED 1 +#endif + +#ifndef SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED +#error "SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED must be defined when the " \ + "builtin is available" +#endif + +// expected-no-diagnostics diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc index 485145734fcbd..7b718bf8c21c5 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc @@ -108,8 +108,11 @@ supports. |Description |1 -|The APIs of this experimental extension are not versioned, so the - feature-test macro always has this value. +|The initial version of this extension. + +|2 +|Adds the `SYCL_EXT_ONEAPI_KERNEL_FUNCTION` macro for launching a free function + kernel by name with template-argument deduction / overload resolution. |=== === Headers @@ -462,6 +465,62 @@ kernel, using the launch configuration specified by `c`. Each value in the `args` pack is passed to the corresponding argument in `Func`, converting it to the argument's type if necessary. +=== Launching a kernel by name + +The launch functions above take the kernel as a `kernel_function` +non-type template argument, which names a fully-resolved function. When the +free function kernel is a function template or an overload set, the application +must therefore name a specific specialization or disambiguate the overload +explicitly (see the example in _Free function kernels which are templates or +overloaded_). + +This extension optionally provides a macro that lets the kernel be named by its +bare identifier, with the specialization resolved from the launch arguments the +same way an ordinary call to the kernel would resolve it: + +``` +SYCL_EXT_ONEAPI_KERNEL_FUNCTION(kernel-name, args...) +``` + +_Effect_: expands to the arguments of an enclosing free function kernel launch +(`single_task` or `nd_launch`). The macro resolves `kernel-name` against the +types of `args` using ordinary C++ overload resolution and template argument +deduction — exactly as the expression `kernel-name(args...)` would — and forms +a launch of the resolved free function kernel with `args` as its kernel +arguments. It is used in the position of the `kernel_function` argument +and the kernel arguments of a launch function, for example: + +``` +// Equivalent to nd_launch(q, ndr, kernel_function>, 3.14f, fptr): +nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(iota, 3.14f, fptr)); + +single_task(q, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(ping, fptr)); +``` + +The `args` are used only to resolve the kernel; they are not evaluated more than +once, and the resolved kernel is launched with those same `args`. + +The return type of the macro is unspecified. Applications must use it only in +the argument position of a free function kernel launch function; the result is +not otherwise usable. + +_Constraints_: The program is ill-formed if `kernel-name`, resolved against the +types of `args`, does not name a SYCL free function kernel (a function declared +with the `nd_range_kernel` or `single_task_kernel` property). + +Ordinary C++ template argument deduction applies to `kernel-name`. A template +parameter of the kernel that cannot be deduced from `args` — for example a +non-type template parameter that does not appear in any function parameter, or a +type parameter used only in a non-deduced context — must be specified +explicitly, just as it would be in a direct call to the kernel (this matches the +behavior of CUDA's `<<<>>>` launch). See the examples below. + +This macro requires support from the device compiler. An implementation that +provides it predefines the macro `SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED` to +`1`; applications can test for this macro to determine whether the feature is +available. The remaining APIs of this extension do not require this support and +are available with any host compiler. + === New kernel bundle member functions This extension adds the following new functions which add kernel bundle support @@ -1117,6 +1176,47 @@ int main() { } ---- +Alternatively, when the implementation defines +`SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED`, the same launches can name the +kernel by its bare identifier and let the compiler deduce the template argument +or resolve the overload from the launch arguments: + +[source,c++] +---- +#ifdef SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED + // T is deduced as float / int from the launch arguments. + syclexp::nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(iota, 3.14f, fptr)); + syclexp::nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(iota, 3, iptr)); + + // The correct ping overload is selected from the argument type. + syclexp::nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(ping, fptr)); + syclexp::nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(ping, iptr)); +#endif +---- + +A template parameter that cannot be deduced from the launch arguments must still +be specified explicitly — this is a property of standard C++ template argument +deduction, not a limitation specific to this macro. For example: + +[source,c++] +---- +// Dim is a non-type template parameter that appears in no function parameter, +// so it cannot be deduced from the launch arguments. +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) +void fill(T *ptr, T value) { /* ... */ } + +#ifdef SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED + // Ill-formed: Dim cannot be deduced (same as writing fill(fptr, 1.0f)). + // syclexp::nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(fill, fptr, 1.0f)); + + // Specify the non-deducible Dim explicitly; T is still deduced. The extra + // parentheses keep the '<...>' comma from being read as a macro-argument + // separator. + syclexp::nd_launch(q, ndr, SYCL_EXT_ONEAPI_KERNEL_FUNCTION((fill<1>), fptr, 1.0f)); +#endif +---- + === Using "scratch" work-group local memory Free function kernels can use work-group local memory via `local_accessor` or diff --git a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp index 81a776eee544a..4a60ab3582b7a 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp @@ -73,3 +73,76 @@ template struct is_device_copyable; } // namespace _V1 } // namespace sycl + +// CUDA-`<<<>>>`-style bare-name launch. SYCL_EXT_ONEAPI_KERNEL_FUNCTION(NAME, +// args...) lets a free function kernel be launched by naming it directly, with +// the compiler deducing its template arguments / resolving the overload from +// the launch arguments, while keeping the existing enqueue-function API +// unchanged: +// +// nd_launch(q, range, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(iota, 3.14f, ptr)); +// single_task(q, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(store42, ptr)); +// +// The macro expands into two comma-separated arguments of the enclosing enqueue +// call: the kernel_function non-type-template selector followed by the +// same launch arguments. __builtin_sycl_launch_kernel resolves NAME against the +// argument types (overload resolution / template argument deduction) and +// evaluates to a pointer to the chosen specialization, which fills the +// `auto *Func` template parameter. The launch arguments are then forwarded +// normally, so the emitted SPIR-V kernel is the real user function (no wrapper) +// and the launch path is exactly kernel_function. +// +// Non-deducible template parameters (those appearing in no function parameter) +// must still be spelled explicitly, e.g. +// SYCL_EXT_ONEAPI_KERNEL_FUNCTION((kern), args); this is a +// fundamental C++ limitation, matching CUDA. +// +// The leading unary '+' on the builtin call is load-bearing, not a typo. When +// SYCL_EXT_ONEAPI_KERNEL_FUNCTION is used inside a template with dependent +// arguments the builtin call is deferred with a BuiltinFn placeholder callee; +// deducing this `auto` non-type template argument classifies the argument +// expression. Wrapping the call in unary '+' makes that argument a +// UnaryOperator (a plain prvalue that is not classified through +// CallExpr::getCallReturnType) instead of the dependent CallExpr, which would +// otherwise assert in the front end. Unary '+' on a function pointer is the +// identity, so the deduced Func value is unchanged; this keeps the whole +// feature free of any compiler-side change for the dependent case. +// +// SYCL_EXT_ONEAPI_KERNEL_FUNCTION is the only part of the free function kernel +// API that needs compiler support: it relies on the +// __builtin_sycl_launch_kernel front-end builtin, which is provided only by the +// SYCL device compiler (Intel oneAPI DPC++ / clang -fsycl). Everything else in +// this extension (SYCL_EXT_ONEAPI_FUNCTION_PROPERTY, kernel_function, +// nd_launch, single_task) is header-only and works with any host compiler +// (MSVC, GCC, +// ...). The macro is therefore gated on __has_builtin: +// * When the builtin is available, SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED +// is defined to 1 and the macro performs the deduction. +// * Otherwise the macro expands to a reference to an undeclared, +// descriptively-named identifier so that *using* it is a clear compile +// error at the call site (portable across GCC / Clang / MSVC), while merely +// including this header remains valid everywhere. +#ifndef __has_builtin +// Older MSVC (< VS2019 16.1) lacks __has_builtin; treat as "not available". +#define __has_builtin(x) 0 +#endif + +#if __has_builtin(__builtin_sycl_launch_kernel) + +#define SYCL_EXT_ONEAPI_KERNEL_FUNCTION_SUPPORTED 1 + +#define SYCL_EXT_ONEAPI_KERNEL_FUNCTION(NAME, ...) \ + ::sycl::ext::oneapi::experimental::kernel_function< \ + +__builtin_sycl_launch_kernel(NAME, ##__VA_ARGS__)>, \ + ##__VA_ARGS__ +#else + +// No compiler support (e.g. a plain MSVC/GCC host compile without the SYCL +// device compiler). Expand to an undeclared identifier whose name IS the +// diagnostic, so any conforming compiler reports a clear "use of undeclared +// identifier" error pointing at the call site. Including the header is fine; +// only using the macro fails. +#define SYCL_EXT_ONEAPI_KERNEL_FUNCTION(NAME, ...) \ + SYCL_EXT_ONEAPI_KERNEL_FUNCTION_requires_dpcpp_as_the_compiler + +#endif diff --git a/sycl/test-e2e/FreeFunctionKernels/sycl_kernel_macro_launch.cpp b/sycl/test-e2e/FreeFunctionKernels/sycl_kernel_macro_launch.cpp new file mode 100644 index 0000000000000..17539dfbbf75e --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/sycl_kernel_macro_launch.cpp @@ -0,0 +1,222 @@ +// REQUIRES: aspect-usm_shared_allocations +// UNSUPPORTED: target-amd +// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16072 + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// XFAIL: target-native_cpu +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20142 + +// CUDA-`<<<>>>`-style bare-name launch via the SYCL_EXT_ONEAPI_KERNEL_FUNCTION +// macro. The existing enqueue-function API (nd_launch / single_task taking a +// kernel_function selector) is kept unchanged; +// SYCL_EXT_ONEAPI_KERNEL_FUNCTION(name, args...) expands to +// kernel_function<__builtin_sycl_launch_kernel(name, args...)>, args... +// so the compiler deduces the kernel's template arguments / resolves the +// overload from the launch arguments while the launched SPIR-V kernel remains +// the real user function (no wrapper). This test exercises single_task and +// nd_launch; non-templated and templated kernels; a dependent (function- +// template) call site; a zero-argument kernel; a kernel with mixed template +// parameters (an explicit non-deducible Dim plus a deduced T); an overload set +// resolved from the launch argument types; and every launch-configuration form +// the macro can target — queue and handler overloads, and the launch_config +// overload. + +#include +#include +#include +#include +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +// Non-templated single_task. +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void store42(int *p) { *p = 42; } + +// Templated single_task. +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void store_one(T *p) { + *p = T{1}; +} + +// Non-templated nd_range. +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void scale(float *y, float k, int n) { + size_t i = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + if (static_cast(i) < n) + y[i] *= k; +} + +// Templated nd_range. +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void axpy(T *y, const T *x, T a, int n) { + size_t i = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + if (static_cast(i) < n) + y[i] = a * x[i] + y[i]; +} + +// Zero-argument single_task kernel (no launch args after the name). +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void tick() {} + +// Mixed template parameters: Dim is a non-deducible NTTP given explicitly at +// the launch site; T is deduced from the launch arguments. The kernel-kind +// property is itself parameterized on Dim. +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) +void fill_val(T *p, T v, int n) { + size_t i = syclext::this_work_item::get_nd_item().get_global_linear_id(); + if (static_cast(i) < n) + p[i] = v; +} + +// Overload set: two free function kernels of the same name distinguished by +// argument type. The builtin resolves which overload from the launch arguments. +// (This required intel/llvm#22793: launching a free function kernel directly +// instead of through a wrapper — the wrapper's integration-header entry could +// not represent an overloaded name.) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void fill(int *p, int v, int n) { + size_t i = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + if (static_cast(i) < n) + p[i] = v; +} +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void fill(float *p, float v, int n) { + size_t i = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + if (static_cast(i) < n) + p[i] = v; +} + +// A SYCL_EXT_ONEAPI_KERNEL_FUNCTION launch from inside an ordinary function +// template, so the launch arguments are dependent and the macro's builtin call +// is deferred to instantiation (regression guard: this dependent-context use +// previously crashed the front end before the builtin call was made +// type-dependent). +template +void run_axpy(sycl::queue q, sycl::nd_range<1> r, T *y, const T *x, T a, + int n) { + syclexp::nd_launch(q, r, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(axpy, y, x, a, n)); +} + +int main() { + sycl::queue q; + constexpr int N = 32; + sycl::nd_range<1> r{sycl::range<1>(N), sycl::range<1>(8)}; + + // single_task, non-templated. + int *p = sycl::malloc_shared(1, q); + *p = 0; + syclexp::single_task(q, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(store42, p)); + q.wait(); + assert(*p == 42); + + // single_task, templated — deduce T = int. + *p = 0; + syclexp::single_task(q, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(store_one, p)); + q.wait(); + assert(*p == 1); + + // nd_range, non-templated. + float *y = sycl::malloc_shared(N, q); + for (int i = 0; i < N; ++i) + y[i] = 2.0f; + syclexp::nd_launch(q, r, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(scale, y, 3.0f, N)); + q.wait(); + for (int i = 0; i < N; ++i) + assert(y[i] == 6.0f); + + // nd_range, templated — deduce T = float. + float *x = sycl::malloc_shared(N, q); + const float a = 2.0f; + for (int i = 0; i < N; ++i) { + x[i] = static_cast(i); + y[i] = 1.0f; + } + syclexp::nd_launch(q, r, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(axpy, y, x, a, N)); + q.wait(); + for (int i = 0; i < N; ++i) + assert(y[i] == a * static_cast(i) + 1.0f); + + // nd_range, templated, from a dependent (function-template) call site: + // deduction is deferred to instantiation. + for (int i = 0; i < N; ++i) + y[i] = 1.0f; + run_axpy(q, r, y, x, a, N); + q.wait(); + for (int i = 0; i < N; ++i) + assert(y[i] == a * static_cast(i) + 1.0f); + + // Zero launch arguments: SYCL_EXT_ONEAPI_KERNEL_FUNCTION(tick) must expand + // without a dangling comma and launch cleanly. + syclexp::single_task(q, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(tick)); + q.wait(); + + // Mixed template params: Dim=1 given explicitly (non-deducible), T deduced + // from the pointer/value arguments. Paren-wrap fill_val<1> so the macro does + // not split on the template-argument comma (harmless here, required in + // general). + for (int i = 0; i < N; ++i) + y[i] = 0.0f; + syclexp::nd_launch( + q, r, SYCL_EXT_ONEAPI_KERNEL_FUNCTION((fill_val<1>), y, 7.0f, N)); + q.wait(); + for (int i = 0; i < N; ++i) + assert(y[i] == 7.0f); + + // Overload set: the builtin selects fill(int*) then fill(float*) purely from + // the launch argument types. + int *ip = sycl::malloc_shared(N, q); + syclexp::nd_launch(q, r, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(fill, ip, 5, N)); + q.wait(); + for (int i = 0; i < N; ++i) + assert(ip[i] == 5); + syclexp::nd_launch(q, r, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(fill, y, 8.0f, N)); + q.wait(); + for (int i = 0; i < N; ++i) + assert(y[i] == 8.0f); + + // The macro is orthogonal to the launch configuration: it also works with the + // handler-taking overloads (inside a command group) and with the + // launch_config overload, not just the queue + nd_range form above. + + // single_task(handler&, ...) inside a command group. + *p = 0; + syclexp::submit(q, [&](sycl::handler &h) { + syclexp::single_task(h, SYCL_EXT_ONEAPI_KERNEL_FUNCTION(store42, p)); + }); + q.wait(); + assert(*p == 42); + + // nd_launch(handler&, nd_range, ...) inside a command group. + for (int i = 0; i < N; ++i) + y[i] = 2.0f; + syclexp::submit(q, [&](sycl::handler &h) { + syclexp::nd_launch(h, r, + SYCL_EXT_ONEAPI_KERNEL_FUNCTION(scale, y, 3.0f, N)); + }); + q.wait(); + for (int i = 0; i < N; ++i) + assert(y[i] == 6.0f); + + // nd_launch(queue, launch_config, ...). + for (int i = 0; i < N; ++i) + y[i] = 2.0f; + syclexp::launch_config> cfg{r}; + syclexp::nd_launch(q, cfg, + SYCL_EXT_ONEAPI_KERNEL_FUNCTION(scale, y, 4.0f, N)); + q.wait(); + for (int i = 0; i < N; ++i) + assert(y[i] == 8.0f); + + sycl::free(p, q); + sycl::free(y, q); + sycl::free(x, q); + sycl::free(ip, q); + return 0; +}