From e1ef61686d636cd8ecbe3c07bd150cc79457749b Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 09:24:03 -0700 Subject: [PATCH 01/16] [KHR][FFK] Add 'Restrictions on kernel argument types' section (device copyable) --- adoc/extensions/index.adoc | 1 + .../sycl_khr_free_function_kernels.adoc | 171 ++++++++++++++++++ 2 files changed, 172 insertions(+) create mode 100644 adoc/extensions/sycl_khr_free_function_kernels.adoc diff --git a/adoc/extensions/index.adoc b/adoc/extensions/index.adoc index d16bc954b..d6a8218ab 100644 --- a/adoc/extensions/index.adoc +++ b/adoc/extensions/index.adoc @@ -18,3 +18,4 @@ include::sycl_khr_queue_flush.adoc[leveloffset=2] include::sycl_khr_work_item_queries.adoc[leveloffset=2] include::sycl_khr_static_addrspace_cast.adoc[leveloffset=2] include::sycl_khr_dynamic_addrspace_cast.adoc[leveloffset=2] +include::sycl_khr_free_function_kernels.adoc[leveloffset=2] diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc new file mode 100644 index 000000000..4d676a0f6 --- /dev/null +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -0,0 +1,171 @@ +[[sec:khr-free-function-kernels]] += sycl_khr_free_function_kernels + +// SKELETON — scaffolding only (Step A + Step B scope trim). Structure + +// boilerplate are real; every content section below is intentionally EMPTY with +// a placeholder comment naming (a) what normative content goes there, (b) which +// section of the experimental spec it ports from +// (sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc), +// and (c) a `// COVERAGE:` marker pointing at the prototype step(s) + test that +// prove it implementable (see playground/ffk-khr-coverage.md for the full +// claim-by-claim map). Do NOT treat any prose here as final. +// +// SCOPE (Step B, project memory "MAJOR SCOPE CUT"): the kernel-bundle and +// native-backend sections are CUT — this KHR is self-contained on core SYCL, +// no vendor/native dependency. Do NOT reintroduce them. The __sycl_kernel_ +// prefix / one-symbol naming is now a dpcpp impl detail + an off-spec opt-in +// flag, NOT a spec concern. + +// INTRO PARAGRAPH (port from experimental "== Overview", :80-110) +// One paragraph, no heading: a free function kernel is an ordinary C++ function, +// decorated so it becomes a device kernel entry point; enables defining kernels +// at namespace scope and launching/looking them up by the function itself. +// Keep it vendor-neutral (no JIT/OpenVINO mention — see project memory charter). +// TODO(Step B): write intro prose. + +[[sec:khr-free-function-kernels-dependencies]] +== Dependencies + +// PORT/NEW: experimental "== Dependencies" (:38-70) lists no extension deps; +// the KHR ADDS two normative dependencies (project memory): +// - sycl_khr_properties (PR #980, gmlueck — UNMERGED forward dep; anchor +// will be sec:khr-properties once landed). Defines the CT/RT property +// type-shape convention (var-template = compile-time, class-w-ctor = +// runtime) that the decoration + launch property surfaces must follow. +// - sycl_khr_work_item_queries (<>) — the in-kernel +// iteration-id queries (this_nd_item etc.) used from a free function body. +// Also note the C++17 baseline (variadic macro is C++17-safe). +// TODO(Step B): write the dependency sentences + cross-refs; flag the #980 +// unmerged-dependency risk. + +[[sec:khr-free-function-kernels-feature-test]] +== Feature test macro + +An implementation supporting this extension must predefine the macro +[code]#SYCL_KHR_FREE_FUNCTION_KERNELS# to one of the values defined in the table +below. + +[%header,cols="1,5"] +|=== +|Value +|Description + +|1 +|Initial version of this extension. +|=== + +[[sec:khr-free-function-kernels-defining]] +== Defining a free function kernel + +// PORT: experimental "=== Defining a free function kernel" (:156-335). +// Define the SYCL_KHR_KERNEL(kind, modifiers...) decoration macro: +// - MANDATORY first arg = kind (single_task_kernel | nd_kernel); empty +// macro is ill-formed (prototype step 12 — empty-macro limbo retired). +// - variadic tail = optional tuning property modifiers (work_group_size, etc.). +// - the function IS the entry point; calling it as an ordinary function or +// from another kernel = undefined behavior. +// - return type must be void; arg constraints -> see arg-types section. +// COVERAGE: step 1 (variadic macro, khr_free_kernel_macro.cpp) + step 12 +// (mandatory kind / empty-is-error, khr_free_kernel_macro.cpp). See coverage map. +// TODO: normative macro definition + decoration semantics. + +[[sec:khr-free-function-kernels-arg-types]] +== Restrictions on kernel argument types + +// COVERAGE: arg type rule = device-copyable, the SAME rule as any SYCL kernel +// argument (<>); reuse core is_device_copyable_v +// rather than a new trait (project memory "ARG-TYPE TWO-TIER DECISION": spec +// tier = DEVICE COPYABLE; the tighter cross-ABI ceiling is off-spec flag-mode +// dpcpp, NOT in this spec). Prototype steps 2/3 (is_valid_kernel_arg_v + +// narrowing) implement the off-spec flag tier, not this section. See +// playground/ffk-khr-coverage.md and ffk-khr-pr-step-C-args-RESULT.md. + +A free function kernel passes each of its kernel arguments as a parameter of the +function, rather than as a member of a function object or a lambda capture. +Each parameter of a free function kernel must have a type that is +<>. +An application can query whether a type [code]#T# is device copyable with the +core [code]#sycl::is_device_copyable_v# type trait. + +A free function kernel admits only device-copyable parameters. +The rules for parameter passing to kernels (<>) +permit two categories of kernel parameter type: any <> type, +and a separate set of special SYCL types — such as [code]#accessor#, +[code]#local_accessor#, the image accessors, [code]#stream#, [code]#reducer#, +and [code]#kernel_handler# — that are legal only because the implementation +passes them in a special, non-positional way. +A free function kernel does not provide this second allowance: only the first +category applies. +A program that uses any of these special types, or any other type that is not +<>, as a free function kernel parameter is ill formed. + +[_Note:_ A free function kernel receives every parameter positionally, as an +ordinary device-copyable kernel argument. +The special types above are not device copyable; they are legal parameters for +general SYCL kernels only through the separate special-passing allowance of +<>, which this extension does not provide. +_{endnote}_] + +[[sec:khr-free-function-kernels-traits]] +== Traits for kernel functions + +// PORT: experimental "=== New traits for kernel functions" (:336-417). +// Compile-time-queryable traits keyed on the kernel function: +// is_kernel_v, is_nd_range_kernel_v, +// is_single_task_kernel_v, plus property queries (has_property / +// get_property). Spec the CONTRACT "queryable at compile time" — these read +// off the host-visible declaration, NO integration header required (project +// memory dissolution chain). Property type-shape follows sycl_khr_properties. +// COVERAGE: step 4 (host-readable __builtin_sycl_has_property/get_property + +// is_kernel/is_nd_kernel/is_single_task_kernel traits; +// sycl-property-builtins.cpp, khr_kernel_property_traits.cpp). Spec promises +// only the OBSERVABLE (traits queryable at compile time), NOT the builtin / +// integration-header mechanism. See coverage map. +// TODO: trait synopses + Constraints/Returns. + +[[sec:khr-free-function-kernels-launch]] +== Launching a free function kernel + +// PORT: experimental "=== New free functions to launch a kernel" (:418-502) + +// "=== Enqueuing a free function kernel and setting parameter values" +// (:800-841). Define khr::nd_launch / khr::single_task (queue + handler forms), +// kernel identity = khr::kernel_function value, both return void. +// Compile-time checks: kind/dimensionality (is_nd_range_kernel_v) + +// arg-type (is_valid_kernel_arg_v) via Constraints. NO compile-time +// work-group-size value check (Greg feedback — constant nd_ranges are rare). +// Runtime must-match (work_group_size etc.) noted as runtime exceptions. +// COVERAGE: step 5 (checked nd_launch/single_task; khr_checked_launch.cpp, +// khr_checked_launch_errors.cpp) + step 7/8 (runtime must-match throws +// errc::nd_range, decoration propagated onto wrapper; e2e khr_must_match.cpp, +// khr_must_match_negative.cpp on PVC — work_group_size proven, max_*/sub_group +// designed-not-e2e-tested). See coverage map. +// TODO: launcher synopses + Constraints/Effects. + +[[sec:khr-free-function-kernels-launch-properties]] +== Launch properties + +// PORT: experimental "=== Interaction with kernel properties" (:864-887). +// Optional khr::properties{} argument to the launchers (Greg: naked property +// list, NOT a launch_config wrapper), carrying RUNTIME-only launch properties. +// Compile-time-value properties are rejected at the launch site (static +// constraint). v1 in-scope payload is minimal (work_group_scratch_size +// deferred); reserve the overload shape. Express which property keys are valid +// at launch via the sycl_khr_properties CT/RT classification. +// COVERAGE: step 6 (runtime-property launch overload, CT props rejected; +// khr_launch_properties.cpp, khr_launch_properties_errors.cpp). v1 in-scope +// payload empty (work_group_scratch_size deferred) — reserve overload shape +// only. See coverage map. +// TODO: overload synopsis + Constraints; defer-list note. + +[[sec:khr-free-function-kernels-example]] +== Example + +// PORT: experimental "== Examples" (:974-1064): "Basic invocation" + +// "Free function kernels which are templates or overloaded". Provide one (or +// two) full compilable programs: #include , define a kernel with +// SYCL_KHR_KERNEL(nd_kernel<1>), launch with khr::nd_launch, verify result. +// Use [source,,linenums]. Mirror the style of +// sycl_khr_work_item_queries.adoc Example. Keep vendor-neutral — NO native +// backend launch in the example (native section cut). +// COVERAGE: steps 5 + 7 (positive e2e, khr_must_match.cpp on PVC). See coverage map. +// TODO: write the example program(s). From 0b6f9233c3970b7d3d01089dec3b7dafc886fb8e Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 09:29:39 -0700 Subject: [PATCH 02/16] [KHR][FFK] Add "Defining a free function kernel" section --- .../sycl_khr_free_function_kernels.adoc | 129 ++++++++++++++++-- 1 file changed, 118 insertions(+), 11 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 4d676a0f6..6ebe81748 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -57,17 +57,124 @@ below. [[sec:khr-free-function-kernels-defining]] == Defining a free function kernel -// PORT: experimental "=== Defining a free function kernel" (:156-335). -// Define the SYCL_KHR_KERNEL(kind, modifiers...) decoration macro: -// - MANDATORY first arg = kind (single_task_kernel | nd_kernel); empty -// macro is ill-formed (prototype step 12 — empty-macro limbo retired). -// - variadic tail = optional tuning property modifiers (work_group_size, etc.). -// - the function IS the entry point; calling it as an ordinary function or -// from another kernel = undefined behavior. -// - return type must be void; arg constraints -> see arg-types section. -// COVERAGE: step 1 (variadic macro, khr_free_kernel_macro.cpp) + step 12 -// (mandatory kind / empty-is-error, khr_free_kernel_macro.cpp). See coverage map. -// TODO: normative macro definition + decoration semantics. +// COVERAGE: step 1 (variadic SYCL_KHR_KERNEL macro, khr_free_kernel_macro.cpp) + +// step 12 (mandatory kind / empty-macro-is-error, khr_free_kernel_macro.cpp). +// Ported from experimental "=== Defining a free function kernel" (:156-205, +// :324-330). Reframed onto the mandatory-kind macro SYCL_KHR_KERNEL(kind, ...) +// (project memory step 12) rather than the experimental +// SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((property)) decoration. Arg-type rule is NOT +// restated here — cross-refs <>. +// SYCL_EXTERNAL is left as an OPEN note (not a normative rule, per project +// memory). Call-as-a-function is stated as UB only (impl is not claimed to +// reject it). See playground/ffk-khr-coverage.md. + +A free function kernel is an ordinary C++ function whose declaration is decorated +with the [code]#SYCL_KHR_KERNEL# macro. +The first argument to the macro is the kernel's _kind_, which is mandatory. +The kind is one of the following: + +* [code]#sycl::khr::single_task_kernel# --- the function is launched as a single + task, without any iteration space; or + +* [code]#sycl::khr::nd_kernel# --- the function is launched over an + [code]#nd_range# iteration space of [code]#Dims# dimensions. + +Any remaining arguments to the macro are optional properties that further +describe the kernel, such as a required work-group size. +These properties are compile-time properties as defined by +<>; this extension does not require any +particular property to be supported. + +[source,role=synopsis,id=api:khr-free-function-kernels-macro] +---- +// Decorate a function declaration to define a free function kernel. +#define SYCL_KHR_KERNEL(kind, ...) /* see below */ +---- + +The first declaration of the function in the translation unit must be the +decorated declaration (see the rules below). +For example: + +[source] +---- +SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +void scale(float factor, float *data) { + // ... +} +---- + +A free function kernel must obey all of the following rules. +A program that violates any of them is ill formed unless stated otherwise. + +* The macro's kind argument must be present. + Invoking the macro with no kind --- [code]#SYCL_KHR_KERNEL()# --- is ill + formed. + +* The function must be declared at namespace scope, or at class scope as a + static member function. + +* The function's return type must be [code]#void#. + +* The function must not accept a variadic argument list. + +* Each of the function's parameters must have a type that is permitted as a free + function kernel argument, as defined in + <>. + +* No declaration of the function may specify a default argument for any + parameter. + +* The decoration must appear on the first declaration of the function in the + translation unit. + A redeclaration of the function may also be decorated, provided it is decorated + with the same kind (and, where present, the same property arguments). + The effect is the same whether or not a redeclaration is decorated. + +* The same function must be decorated with exactly one kind. + Decorating the same function with more than one kind, or with the same kind but + conflicting property arguments, is ill formed. + Decorating it more than once with the same kind and the same property arguments + is permitted and has the same effect as decorating it once. + +* If the function is decorated in one translation unit, every other translation + unit that declares the same function must decorate it identically (the same + kind and the same property arguments). + +* The way the function's body obtains its iteration ID must be consistent with + the declared kind: a function whose kind is + [code]#sycl::khr::nd_kernel# must obtain its iteration ID for an + [code]#nd_range# of [code]#Dims# dimensions, and a function whose kind is + [code]#sycl::khr::single_task_kernel# must not obtain an iteration ID. + A program that violates this rule is ill formed, no diagnostic required. + +[_Note:_ An implementation may, but need not, diagnose an inconsistency between +the function's body and its declared kind. +Such a diagnosis is not always possible, for example when the body is provided in +another translation unit or as a precompiled device object. +_{endnote}_] + +A function decorated with [code]#SYCL_KHR_KERNEL# is a kernel entry point. +Calling it as an ordinary function, in either host or device code, results in +undefined behavior. +The function's address may still be taken; this address identifies the kernel and +is the value passed to the launch and query APIs described below (the +[code]#Func# parameter). + +[_Note:_ Because calling a free function kernel as an ordinary function is +undefined behavior rather than an ill-formed program, an implementation is not +required to diagnose such a call. +An implementation is nevertheless permitted, as a quality-of-implementation +matter, to diagnose a call that appears in host code. +_{endnote}_] + +// OPEN (raise in PR discussion, not a normative rule): may a free function +// kernel entry point be declared SYCL_EXTERNAL? Tension — the host needs the +// symbol visible (to name/launch it) but never calls it, while SYCL_EXTERNAL is +// a device cross-TU linkage contract that is meaningless for an entry point +// (calling a kernel from device code is UB). The experimental spec (:204) +// contemplates SYCL_EXTERNAL free function kernel bodies; whether the KHR should +// permit, forbid, or stay silent on this is deferred to the PR. Do NOT add a +// normative SYCL_EXTERNAL rule here until that is resolved. [[sec:khr-free-function-kernels-arg-types]] == Restrictions on kernel argument types From 3823195a96235d98f1165502be2bc403540940de Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 09:35:35 -0700 Subject: [PATCH 03/16] [KHR][FFK] Add "Traits for kernel functions" section --- .../sycl_khr_free_function_kernels.adoc | 111 ++++++++++++++++-- 1 file changed, 98 insertions(+), 13 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 6ebe81748..fa6086e21 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -216,19 +216,104 @@ _{endnote}_] [[sec:khr-free-function-kernels-traits]] == Traits for kernel functions -// PORT: experimental "=== New traits for kernel functions" (:336-417). -// Compile-time-queryable traits keyed on the kernel function: -// is_kernel_v, is_nd_range_kernel_v, -// is_single_task_kernel_v, plus property queries (has_property / -// get_property). Spec the CONTRACT "queryable at compile time" — these read -// off the host-visible declaration, NO integration header required (project -// memory dissolution chain). Property type-shape follows sycl_khr_properties. -// COVERAGE: step 4 (host-readable __builtin_sycl_has_property/get_property + -// is_kernel/is_nd_kernel/is_single_task_kernel traits; -// sycl-property-builtins.cpp, khr_kernel_property_traits.cpp). Spec promises -// only the OBSERVABLE (traits queryable at compile time), NOT the builtin / -// integration-header mechanism. See coverage map. -// TODO: trait synopses + Constraints/Returns. +// COVERAGE: step 4 (is_kernel / is_nd_kernel / is_single_task_kernel traits +// queryable at host compile time off the decorated declaration, NO integration +// header — sycl-property-builtins.cpp, khr_kernel_property_traits.cpp). Ported +// from experimental "=== New traits for kernel functions" (:336-417); renamed +// is_nd_range_kernel -> is_nd_kernel (khr kind name) and moved to sycl::khr. +// Spec states ONLY the OBSERVABLE contract ("usable in constant expressions / +// SFINAE at compile time"); the builtin / integration-header / device-compile +// mechanism is an implementation detail and is deliberately not specified. +// A general has_property / get_property query surface is DEFERRED (it belongs +// with sycl_khr_properties, #980) — only the three kind/dimensionality traits +// are exposed here, matching the experimental spec. See coverage map. + +This extension defines traits that report, at compile time, whether a given +[code]#Func# is a free function kernel (see +<>) and, if so, what kind it has. +In each trait, [code]#Func# is the address of a function; the trait reports +[code]#true# only when [code]#Func# is the address of a function that is +decorated as a free function kernel of the relevant kind. + +These traits are usable in constant expressions and in SFINAE contexts at +compile time, so an application (or a software layer built on top of this +extension) can branch on, or constrain a template against, the kind of a +free function kernel without launching it. + +{note}The kernel whose address is [code]#Func# need not be defined in the same +translation unit as the use of the trait; the declaration decorated with +[code]#SYCL_KHR_KERNEL# is sufficient. +How an implementation makes the decoration observable in a constant expression +is unspecified.{endnote} + +''' + +.[apidef]#is_kernel# +[source,role=synopsis,id=api:khr-free-function-kernels-is_kernel] +---- +namespace sycl::khr { + +template +struct is_kernel; + +template +inline constexpr bool is_kernel_v = is_kernel::value; + +} // namespace sycl::khr +---- + +_Returns:_ [code]#is_kernel::value# is [code]#true# if [code]#Func# is the +address of a function that is decorated as a free function kernel (with any +kind), and [code]#false# otherwise. + +_Remarks:_ The expression [code]#is_kernel_v# is usable in a constant +expression. + +''' + +.[apidef]#is_nd_kernel# +[source,role=synopsis,id=api:khr-free-function-kernels-is_nd_kernel] +---- +namespace sycl::khr { + +template +struct is_nd_kernel; + +template +inline constexpr bool is_nd_kernel_v = is_nd_kernel::value; + +} // namespace sycl::khr +---- + +_Returns:_ [code]#is_nd_kernel::value# is [code]#true# if +[code]#Func# is the address of a function whose kind is +[code]#sycl::khr::nd_kernel#, and [code]#false# otherwise. + +_Remarks:_ The expression [code]#is_nd_kernel_v# is usable in a +constant expression. + +''' + +.[apidef]#is_single_task_kernel# +[source,role=synopsis,id=api:khr-free-function-kernels-is_single_task_kernel] +---- +namespace sycl::khr { + +template +struct is_single_task_kernel; + +template +inline constexpr bool is_single_task_kernel_v = is_single_task_kernel::value; + +} // namespace sycl::khr +---- + +_Returns:_ [code]#is_single_task_kernel::value# is [code]#true# if +[code]#Func# is the address of a function whose kind is +[code]#sycl::khr::single_task_kernel#, and [code]#false# otherwise. + +_Remarks:_ The expression [code]#is_single_task_kernel_v# is usable in a +constant expression. [[sec:khr-free-function-kernels-launch]] == Launching a free function kernel From bb49fd9c25b8b714bbfe6e3c192088690aa05847 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 11:02:33 -0700 Subject: [PATCH 04/16] [KHR][FFK] Add "Launching a free function kernel" section --- .../sycl_khr_free_function_kernels.adoc | 124 +++++++++++++++++- 1 file changed, 123 insertions(+), 1 deletion(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index fa6086e21..67397f4d9 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -331,7 +331,129 @@ constant expression. // errc::nd_range, decoration propagated onto wrapper; e2e khr_must_match.cpp, // khr_must_match_negative.cpp on PVC — work_group_size proven, max_*/sub_group // designed-not-e2e-tested). See coverage map. -// TODO: launcher synopses + Constraints/Effects. + +A free function kernel is launched by passing its address to one of the launch +functions defined in this section. +The kernel's address is supplied through the [code]#kernel_function# handle, +which carries the address as a template parameter. +A launch function deduces the kernel function [code]#Func# from this handle, so +the kernel is identified by the _value_ [code]#sycl::khr::kernel_function# +that is passed as an argument --- not by an explicit template argument on the +launch function itself. + +.[apidef]#kernel_function# +[source,role=synopsis,id=api:khr-free-function-kernels-kernel_function] +---- +namespace sycl::khr { + +template +struct kernel_function_s {}; + +template +inline constexpr kernel_function_s kernel_function; + +} // namespace sycl::khr +---- + +_Remarks:_ [code]#Func# is the address of a function that is decorated as a free +function kernel (see <>). +The variable template [code]#kernel_function# is the value passed to +[code]#single_task# and [code]#nd_launch# to identify the kernel to launch. + +{note}For a free function kernel that is a template, the address of a particular +specialization is used, for example +[code]#kernel_function>#.{endnote} + +[[sec:khr-free-function-kernels-launch-single_task]] +=== single_task + +''' + +.[apidef]#single_task# +[source,role=synopsis,id=api:khr-free-function-kernels-single_task] +---- +namespace sycl::khr { + +template +void single_task(queue q, kernel_function_s k, Args&&... args); + +template +void single_task(handler &h, kernel_function_s k, Args&&... args); + +} // namespace sycl::khr +---- + +_Constraints:_ Available only if [code]#is_single_task_kernel_v# is +[code]#true# and [code]#std::is_invocable_v# is +[code]#true#. + +_Effects:_ Enqueues the free function kernel [code]#Func# to the queue +[code]#q# (or to the command group of the handler [code]#h#) as a single task. +Each value in the [code]#args# pack is passed to the corresponding parameter of +[code]#Func#, converting it to the parameter's type if necessary. + +[[sec:khr-free-function-kernels-launch-nd_launch]] +=== nd_launch + +''' + +.[apidef]#nd_launch# +[source,role=synopsis,id=api:khr-free-function-kernels-nd_launch] +---- +namespace sycl::khr { + +template +void nd_launch(queue q, nd_range r, + kernel_function_s k, Args&&... args); + +template +void nd_launch(handler &h, nd_range r, + kernel_function_s k, Args&&... args); + +} // namespace sycl::khr +---- + +_Constraints:_ Available only if [code]#is_nd_kernel_v# is +[code]#true# and [code]#std::is_invocable_v# is +[code]#true#. + +_Effects:_ Enqueues the free function kernel [code]#Func# to the queue +[code]#q# (or to the command group of the handler [code]#h#) as an ND-range +kernel, using the number of work-items specified by the [code]#nd_range# +[code]#r#. +Each value in the [code]#args# pack is passed to the corresponding parameter of +[code]#Func#, converting it to the parameter's type if necessary. + +The launch functions enqueue the kernel by its address; the kernel is selected +by the [code]#kernel_function# value, as in: + +[source] +---- +sycl::khr::nd_launch(q, r, sycl::khr::kernel_function, factor, data); +---- + +The [code]#Constraints# above check, at compile time, only the kernel's _kind_ +and _dimensionality_ and that the supplied arguments can be passed to +[code]#Func#. +They do not check the values carried by the kernel's compile-time properties, +such as a required work-group size: a required work-group size is compared +against the local size of the [code]#nd_range# at launch, and the local size of +an [code]#nd_range# is generally a run-time value, so this comparison is +performed at run time rather than at compile time. + +Any launch requirements that a free function kernel carries through its +compile-time properties are enforced at launch in the same way as for any other +SYCL kernel, as described in <>. +In particular, launching a kernel that declares a required work-group size with +an [code]#nd_range# whose local size does not match that requirement throws a +synchronous [code]#exception# with the [code]#errc::nd_range# error code, +exactly as for a SYCL kernel decorated with [code]#reqd_work_group_size#. + +{note}A future version of this extension is expected to add launch functions +that accept a launch property list, allowing launch-time properties --- such as +a dynamic work-group local memory size --- to be specified when a kernel is +launched. +See <>.{endnote} [[sec:khr-free-function-kernels-launch-properties]] == Launch properties From cb33535b1649f6cb68ffbb66535dc1180a4ed6ec Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 13:33:28 -0700 Subject: [PATCH 05/16] [KHR][FFK] Add "Kernel properties" section (free_function_kernel property tag) --- .../sycl_khr_free_function_kernels.adoc | 115 +++++++++++++++++- 1 file changed, 110 insertions(+), 5 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 67397f4d9..518e20417 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -79,11 +79,12 @@ The kind is one of the following: * [code]#sycl::khr::nd_kernel# --- the function is launched over an [code]#nd_range# iteration space of [code]#Dims# dimensions. -Any remaining arguments to the macro are optional properties that further -describe the kernel, such as a required work-group size. -These properties are compile-time properties as defined by -<>; this extension does not require any -particular property to be supported. +Any remaining arguments to the macro are optional compile-time properties that +are applicable to a free function kernel, as described in +<>. +Apart from the _kind_ properties, this extension does not define any such +property; any additional property is defined by the SYCL implementation or by +another extension. [source,role=synopsis,id=api:khr-free-function-kernels-macro] ---- @@ -176,6 +177,110 @@ _{endnote}_] // permit, forbid, or stay silent on this is deferred to the PR. Do NOT add a // normative SYCL_EXTERNAL rule here until that is resolved. +[[sec:khr-free-function-kernels-properties]] +== Kernel properties + +// COVERAGE: the _kind_ properties (single_task_kernel / nd_kernel) are +// PROVEN (step 1 macro + step 12 mandatory-kind validator, +// khr_free_kernel_macro.cpp). The free_function_kernel TAG TYPE and the +// is_property_for-keyed applicability model are a NEW SPEC CONSTRUCT (design +// proposal, project memory "FFK NEEDS A PROPERTY-TAG TYPE"): the prototype used +// khr::detail introspection of the property bundle, NOT a public tag. The +// compile-time vs runtime split reuses Greg's is_property_key_compile_time +// discriminator (step 6, khr_launch_properties*.cpp). DEPENDS on the unmerged +// sycl_khr_properties (#980) — the is_property_for / is_property_key_compile_time +// cross-refs below dangle (sec:khr-properties) until it lands. See +// playground/ffk-khr-coverage.md and ffk-khr-pr-step-C5-property-model.md. + +A free function kernel can carry _properties_, using the property infrastructure +defined by <>. +There are two contexts in which a property is associated with a free function +kernel: + +* _Compile-time properties_ are placed on the kernel by listing them as the + remaining arguments of the [code]#SYCL_KHR_KERNEL# decoration (see + <>). + +* _Runtime properties_ are supplied when the kernel is launched, through the + launch property list (see <>). + +[[sec:khr-free-function-kernels-properties-tag]] +=== The free_function_kernel property tag + +The [code]#sycl_khr_properties# extension expresses the set of classes that a +property may be used with through the [api]#khr::is_property_for# trait, which is +keyed on a _class_ type: [code]#khr::is_property_for_v# is [code]#true# +when the property [code]#P# can be used with the class [code]#Class#. +A free function kernel, however, is a decorated _function_ rather than an object +of some class, so there is no class type to use as the [code]#Class# argument. +This extension therefore introduces an empty tag type that stands in for the +free function kernel as that class. + +''' + +.[apidef]#khr::free_function_kernel# +[source,role=synopsis,id=api:khr-free-function-kernels-tag] +---- +namespace sycl::khr { + +struct free_function_kernel; + +} // namespace sycl::khr +---- + +_Remarks:_ [code]#free_function_kernel# is an incomplete tag type that is never +instantiated. +It serves only as the [code]#Class# argument of the +<> property traits, identifying the +properties that may be associated with a free function kernel. + +A property [code]#P# is _applicable to a free function kernel_ if, and only if, +[code]#khr::is_property_for_v# is [code]#true#. + +{note}A tag type is needed because [code]#khr::is_property_for# keys property +applicability on a class, and a free function kernel has no class object to key +on. +The [code]#free_function_kernel# tag stands in for that class, so the existing +<> machinery is reused unchanged: no new +applicability trait is introduced.{endnote} + +[[sec:khr-free-function-kernels-properties-ct-rt]] +=== Compile-time and runtime properties + +Whether a property is associated with a free function kernel through the +decoration or through the launch property list is determined by whether it is a +compile-time property or a runtime property, as classified by the +[api]#khr::is_property_key_compile_time# trait of +<>. + +* The [code]#SYCL_KHR_KERNEL# decoration accepts a property [code]#P# only if + [code]#P# is applicable to a free function kernel and the key of [code]#P# is a + compile-time property, that is + [code]#khr::is_property_key_compile_time_v# is [code]#true#. + +* The launch property list (see + <>) accepts a property + [code]#P# only if [code]#P# is applicable to a free function kernel and the key + of [code]#P# is a runtime property, that is + [code]#khr::is_property_key_compile_time_v# is [code]#false#. + +[[sec:khr-free-function-kernels-properties-kind]] +=== Properties defined by this extension + +This extension defines only the _kind_ properties, +[code]#sycl::khr::single_task_kernel# and [code]#sycl::khr::nd_kernel# +(see <> for their meaning). +These are compile-time properties that are applicable to a free function kernel, +so they fit the model described above and may appear as the mandatory first +argument of the [code]#SYCL_KHR_KERNEL# decoration. + +As with all SYCL properties, properties can only be defined by the SYCL +implementation, so the SYCL specification (or an extension specification) +provides a formal specification for each property. +Apart from the _kind_ properties, any property that is applicable to a free +function kernel is defined by the SYCL implementation or by another extension +specification; this extension does not define any such property. + [[sec:khr-free-function-kernels-arg-types]] == Restrictions on kernel argument types From 0815c61d56efee3bd0901a19a7d47b45810961e4 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 13:41:59 -0700 Subject: [PATCH 06/16] [KHR][FFK] Add "Launch properties" section (runtime-property launch overloads) --- .../sycl_khr_free_function_kernels.adoc | 142 ++++++++++++++++-- 1 file changed, 126 insertions(+), 16 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 518e20417..f93c6c56d 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -554,27 +554,137 @@ an [code]#nd_range# whose local size does not match that requirement throws a synchronous [code]#exception# with the [code]#errc::nd_range# error code, exactly as for a SYCL kernel decorated with [code]#reqd_work_group_size#. -{note}A future version of this extension is expected to add launch functions -that accept a launch property list, allowing launch-time properties --- such as -a dynamic work-group local memory size --- to be specified when a kernel is -launched. -See <>.{endnote} +{note}The launch functions above do not accept a launch property list. +Overloads that do are defined in +<>; this version of the +extension reserves the shape of those overloads but defines no runtime property +that can be passed through them.{endnote} [[sec:khr-free-function-kernels-launch-properties]] == Launch properties -// PORT: experimental "=== Interaction with kernel properties" (:864-887). -// Optional khr::properties{} argument to the launchers (Greg: naked property -// list, NOT a launch_config wrapper), carrying RUNTIME-only launch properties. -// Compile-time-value properties are rejected at the launch site (static -// constraint). v1 in-scope payload is minimal (work_group_scratch_size -// deferred); reserve the overload shape. Express which property keys are valid -// at launch via the sycl_khr_properties CT/RT classification. // COVERAGE: step 6 (runtime-property launch overload, CT props rejected; -// khr_launch_properties.cpp, khr_launch_properties_errors.cpp). v1 in-scope -// payload empty (work_group_scratch_size deferred) — reserve overload shape -// only. See coverage map. -// TODO: overload synopsis + Constraints; defer-list note. +// khr_launch_properties.cpp, khr_launch_properties_errors.cpp). The overload +// shape is RESERVED on BOTH nd_launch and single_task (project memory +// "LAUNCH-PROPERTY SECTION DESIGN"): symmetric and future-proof, even though +// single_task has no in-scope launch property in this version. The CT/RT split +// reuses Greg's is_property_key_compile_time discriminator; applicability reuses +// the free_function_kernel tag (C5). v1 in-scope payload is EMPTY +// (work_group_scratch_size is deferred to another extension) — this version +// defines no runtime launch property. DEPENDS on the unmerged sycl_khr_properties +// (#980): is_property_list_for / is_property_key_compile_time / sec:khr-properties +// dangle until it lands. See playground/ffk-khr-coverage.md. + +The launch functions in <> have +additional overloads that accept a _launch property list_ --- a +<> properties list supplied at the launch +site. +The property list is passed as a separate argument, immediately before the +[code]#kernel_function# handle; it is not wrapped in any launch-configuration +type. +The properties in the list carry _runtime_ information that applies to the +individual launch, as distinct from the compile-time properties that decorate the +kernel itself (see <>). + +A launch property list accepts only properties that are _applicable to a free +function kernel_ (see <>) and whose +keys are _runtime properties_. +This version of the extension defines no such property: the in-scope set of +runtime launch properties is empty, and the overloads below exist to reserve the +shape that future runtime launch properties will use. + +{note}A runtime launch property is one whose value is not known until the kernel +is launched, so it cannot be carried by the compile-time decoration on the kernel. +A likely future example, defined by another extension, is a dynamic work-group +local memory size, whose byte count is a run-time value supplied at launch. +This extension reserves the launch-property overloads so that such a property can +be added without changing the shape of the launch API.{endnote} + +[[sec:khr-free-function-kernels-launch-properties-single_task]] +=== single_task + +''' + +.[apidef]#single_task# +[source,role=synopsis,id=api:khr-free-function-kernels-single_task-props] +---- +namespace sycl::khr { + +template +void single_task(queue q, Properties props, + kernel_function_s k, Args&&... args); + +template +void single_task(handler &h, Properties props, + kernel_function_s k, Args&&... args); + +} // namespace sycl::khr +---- + +_Constraints:_ Available only if all of the following hold: + +* [code]#Properties# is a <> properties + list, that is [code]#khr::is_property_list_for_v# + is [code]#true#. + This requires every property in [code]#props# to be applicable to a free + function kernel (see <>). +* For every property in [code]#props#, the key of that property is a runtime + property, that is [code]#khr::is_property_key_compile_time_v# is + [code]#false#. + Supplying a compile-time property in the launch property list is ill formed. +* [code]#is_single_task_kernel_v# is [code]#true# and + [code]#std::is_invocable_v# is [code]#true#, as for + the launch function without a property list. + +_Effects:_ Equivalent to the [code]#single_task# overload without a property list +(see <>), additionally applying +the runtime properties in [code]#props# to the launch. + +[[sec:khr-free-function-kernels-launch-properties-nd_launch]] +=== nd_launch + +''' + +.[apidef]#nd_launch# +[source,role=synopsis,id=api:khr-free-function-kernels-nd_launch-props] +---- +namespace sycl::khr { + +template +void nd_launch(queue q, nd_range r, Properties props, + kernel_function_s k, Args&&... args); + +template +void nd_launch(handler &h, nd_range r, Properties props, + kernel_function_s k, Args&&... args); + +} // namespace sycl::khr +---- + +_Constraints:_ Available only if all of the following hold: + +* [code]#Properties# is a <> properties + list, that is [code]#khr::is_property_list_for_v# + is [code]#true#. + This requires every property in [code]#props# to be applicable to a free + function kernel (see <>). +* For every property in [code]#props#, the key of that property is a runtime + property, that is [code]#khr::is_property_key_compile_time_v# is + [code]#false#. + Supplying a compile-time property in the launch property list is ill formed. +* [code]#is_nd_kernel_v# is [code]#true# and + [code]#std::is_invocable_v# is [code]#true#, as for + the launch function without a property list. + +_Effects:_ Equivalent to the [code]#nd_launch# overload without a property list +(see <>), additionally applying +the runtime properties in [code]#props# to the launch. + +{note}Because this version of the extension defines no runtime property that is +applicable to a free function kernel, the only property list that satisfies the +[code]#Constraints# above is the empty properties list. +The overloads are nonetheless specified so that the launch API does not change +shape when a future extension adds a runtime launch property.{endnote} [[sec:khr-free-function-kernels-example]] == Example From 0a3463bfb0213450e01ef6c3fe00a97c6b5ec4a6 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 13:48:59 -0700 Subject: [PATCH 07/16] [KHR][FFK] Add "Dependencies" section --- .../sycl_khr_free_function_kernels.adoc | 38 +++++++++++++------ 1 file changed, 27 insertions(+), 11 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index f93c6c56d..f497e484e 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -26,17 +26,33 @@ [[sec:khr-free-function-kernels-dependencies]] == Dependencies -// PORT/NEW: experimental "== Dependencies" (:38-70) lists no extension deps; -// the KHR ADDS two normative dependencies (project memory): -// - sycl_khr_properties (PR #980, gmlueck — UNMERGED forward dep; anchor -// will be sec:khr-properties once landed). Defines the CT/RT property -// type-shape convention (var-template = compile-time, class-w-ctor = -// runtime) that the decoration + launch property surfaces must follow. -// - sycl_khr_work_item_queries (<>) — the in-kernel -// iteration-id queries (this_nd_item etc.) used from a free function body. -// Also note the C++17 baseline (variadic macro is C++17-safe). -// TODO(Step B): write the dependency sentences + cross-refs; flag the #980 -// unmerged-dependency risk. +// COVERAGE: this section OWNS the #980 (sycl_khr_properties) unmerged +// forward-dependency caveat for the whole document — the editorial note below is +// the single place that explains why every <> cross-ref +// in this extension dangles in a build until #980 lands. Do not restate the +// property model here (see <>). See +// playground/ffk-khr-pr-step-C7-dependencies.md. + +This extension depends on the +<> extension, whose property +infrastructure it uses to associate properties with a free function kernel --- +both the compile-time properties that decorate the kernel and the runtime +properties supplied through the launch property list (see +<>). + +This extension depends on the +<> extension, which +provides the in-kernel queries (such as [code]#sycl::khr::this_nd_item#) that the +body of a [code]#sycl::khr::nd_kernel# free function kernel uses to obtain +its work-item's position in the iteration space. + +{note}At the time of writing, [code]#sycl_khr_properties# is a proposed extension +(SYCL specification PR #980) and is not yet part of the SYCL specification. +Until it is incorporated, the [code]#sec:khr-properties# cross-references in this +extension do not resolve.{endnote} + +Some features of this extension are only available when a SYCL implementation +conforms to {cpp17} or later. [[sec:khr-free-function-kernels-feature-test]] == Feature test macro From 12ccab827da52864e93b17eac85195a5e2ed426b Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 13:54:41 -0700 Subject: [PATCH 08/16] [KHR][FFK] Add overview intro and example --- .../sycl_khr_free_function_kernels.adoc | 119 +++++++++++++----- 1 file changed, 89 insertions(+), 30 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index f497e484e..eb61f5dba 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -1,27 +1,16 @@ [[sec:khr-free-function-kernels]] = sycl_khr_free_function_kernels -// SKELETON — scaffolding only (Step A + Step B scope trim). Structure + -// boilerplate are real; every content section below is intentionally EMPTY with -// a placeholder comment naming (a) what normative content goes there, (b) which -// section of the experimental spec it ports from -// (sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc), -// and (c) a `// COVERAGE:` marker pointing at the prototype step(s) + test that -// prove it implementable (see playground/ffk-khr-coverage.md for the full -// claim-by-claim map). Do NOT treat any prose here as final. -// -// SCOPE (Step B, project memory "MAJOR SCOPE CUT"): the kernel-bundle and -// native-backend sections are CUT — this KHR is self-contained on core SYCL, -// no vendor/native dependency. Do NOT reintroduce them. The __sycl_kernel_ -// prefix / one-symbol naming is now a dpcpp impl detail + an off-spec opt-in -// flag, NOT a spec concern. - -// INTRO PARAGRAPH (port from experimental "== Overview", :80-110) -// One paragraph, no heading: a free function kernel is an ordinary C++ function, -// decorated so it becomes a device kernel entry point; enables defining kernels -// at namespace scope and launching/looking them up by the function itself. -// Keep it vendor-neutral (no JIT/OpenVINO mention — see project memory charter). -// TODO(Step B): write intro prose. +This extension introduces _free function kernels_. +A free function kernel is an ordinary C++ function, defined at namespace scope or +as a static member function, that is decorated with the [code]#SYCL_KHR_KERNEL# +macro so that it becomes a device kernel entry point. +Its kernel arguments are the parameters of the function, rather than the captures +of a lambda expression or the member variables of a function object, which gives +the arguments a defined order. +Because the kernel is a named function, it is identified by the function itself +--- by its address --- so it can be launched and queried by reference to that +function, instead of only as a lambda. [[sec:khr-free-function-kernels-dependencies]] == Dependencies @@ -705,12 +694,82 @@ shape when a future extension adds a runtime launch property.{endnote} [[sec:khr-free-function-kernels-example]] == Example -// PORT: experimental "== Examples" (:974-1064): "Basic invocation" + -// "Free function kernels which are templates or overloaded". Provide one (or -// two) full compilable programs: #include , define a kernel with -// SYCL_KHR_KERNEL(nd_kernel<1>), launch with khr::nd_launch, verify result. -// Use [source,,linenums]. Mirror the style of -// sycl_khr_work_item_queries.adoc Example. Keep vendor-neutral — NO native -// backend launch in the example (native section cut). -// COVERAGE: steps 5 + 7 (positive e2e, khr_must_match.cpp on PVC). See coverage map. -// TODO: write the example program(s). +// COVERAGE: steps 5 + 7 (positive e2e, khr_must_match.cpp on PVC). See +// playground/ffk-khr-coverage.md. + +The example below defines a free function kernel at namespace scope, decorates it +with [code]#SYCL_KHR_KERNEL#, and launches it with [code]#sycl::khr::nd_launch#. +The kernel obtains its work-item position with the +<> query +[code]#sycl::khr::this_nd_item#, rather than receiving an [code]#nd_item# +parameter. + +[source,,linenums] +---- +#include + +constexpr size_t N = 1024; +constexpr size_t WGSIZE = 32; + +// A free function kernel: an ordinary function, decorated with the kind. +SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +void scale(float factor, float *data) { + size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); + data[i] *= factor; +} + +int main() { + sycl::queue q; + + float *data = sycl::malloc_shared(N, q); + for (size_t i = 0; i < N; ++i) + data[i] = static_cast(i); + + // Identify the kernel by the function itself and launch it. + sycl::khr::nd_launch(q, sycl::nd_range<1>{N, WGSIZE}, + sycl::khr::kernel_function, 2.0f, data); + q.wait(); + + for (size_t i = 0; i < N; ++i) + assert(data[i] == 2.0f * static_cast(i)); + + sycl::free(data, q); + return 0; +} +---- + +A free function kernel may also be a function template. +The address passed through [code]#kernel_function# is then the address of a +particular specialization. + +[source,,linenums] +---- +#include + +constexpr size_t N = 1024; +constexpr size_t WGSIZE = 32; + +template +SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +void fill(T value, T *data) { + size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); + data[i] = value; +} + +int main() { + sycl::queue q; + + int *data = sycl::malloc_shared(N, q); + + // For a templated kernel, pass the address of a specific instantiation. + sycl::khr::nd_launch(q, sycl::nd_range<1>{N, WGSIZE}, + sycl::khr::kernel_function>, 42, data); + q.wait(); + + for (size_t i = 0; i < N; ++i) + assert(data[i] == 42); + + sycl::free(data, q); + return 0; +} +---- From e12783519765b5e9e5e81bb0b205830a565966bb Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 14:11:18 -0700 Subject: [PATCH 09/16] [KHR][FFK] Tighten spec prose to KHR density (de-dup + concision) --- .../sycl_khr_free_function_kernels.adoc | 135 +++++++----------- 1 file changed, 48 insertions(+), 87 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index eb61f5dba..8158fdf27 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -5,12 +5,10 @@ This extension introduces _free function kernels_. A free function kernel is an ordinary C++ function, defined at namespace scope or as a static member function, that is decorated with the [code]#SYCL_KHR_KERNEL# macro so that it becomes a device kernel entry point. -Its kernel arguments are the parameters of the function, rather than the captures -of a lambda expression or the member variables of a function object, which gives -the arguments a defined order. -Because the kernel is a named function, it is identified by the function itself ---- by its address --- so it can be launched and queried by reference to that -function, instead of only as a lambda. +Its kernel arguments are the function's parameters, rather than the captures of a +lambda or the members of a function object, which gives them a defined order. +Because it is a named function, the kernel is identified by its address, so it can +be launched and queried by reference to the function rather than only as a lambda. [[sec:khr-free-function-kernels-dependencies]] == Dependencies @@ -24,9 +22,7 @@ function, instead of only as a lambda. This extension depends on the <> extension, whose property -infrastructure it uses to associate properties with a free function kernel --- -both the compile-time properties that decorate the kernel and the runtime -properties supplied through the launch property list (see +infrastructure it uses to associate properties with a free function kernel (see <>). This extension depends on the @@ -37,8 +33,8 @@ its work-item's position in the iteration space. {note}At the time of writing, [code]#sycl_khr_properties# is a proposed extension (SYCL specification PR #980) and is not yet part of the SYCL specification. -Until it is incorporated, the [code]#sec:khr-properties# cross-references in this -extension do not resolve.{endnote} +Until it is incorporated, the cross-references to it in this extension do not +resolve.{endnote} Some features of this extension are only available when a SYCL implementation conforms to {cpp17} or later. @@ -87,9 +83,6 @@ The kind is one of the following: Any remaining arguments to the macro are optional compile-time properties that are applicable to a free function kernel, as described in <>. -Apart from the _kind_ properties, this extension does not define any such -property; any additional property is defined by the SYCL implementation or by -another extension. [source,role=synopsis,id=api:khr-free-function-kernels-macro] ---- @@ -166,11 +159,9 @@ The function's address may still be taken; this address identifies the kernel an is the value passed to the launch and query APIs described below (the [code]#Func# parameter). -[_Note:_ Because calling a free function kernel as an ordinary function is -undefined behavior rather than an ill-formed program, an implementation is not -required to diagnose such a call. -An implementation is nevertheless permitted, as a quality-of-implementation -matter, to diagnose a call that appears in host code. +[_Note:_ Because such a call is undefined behavior rather than ill formed, an +implementation is not required to diagnose it, though it is permitted to diagnose +one that appears in host code as a quality-of-implementation matter. _{endnote}_] // OPEN (raise in PR discussion, not a normative rule): may a free function @@ -219,7 +210,9 @@ when the property [code]#P# can be used with the class [code]#Class#. A free function kernel, however, is a decorated _function_ rather than an object of some class, so there is no class type to use as the [code]#Class# argument. This extension therefore introduces an empty tag type that stands in for the -free function kernel as that class. +free function kernel as that class, so the existing +<> machinery is reused unchanged with no +new applicability trait. ''' @@ -242,13 +235,6 @@ properties that may be associated with a free function kernel. A property [code]#P# is _applicable to a free function kernel_ if, and only if, [code]#khr::is_property_for_v# is [code]#true#. -{note}A tag type is needed because [code]#khr::is_property_for# keys property -applicability on a class, and a free function kernel has no class object to key -on. -The [code]#free_function_kernel# tag stands in for that class, so the existing -<> machinery is reused unchanged: no new -applicability trait is introduced.{endnote} - [[sec:khr-free-function-kernels-properties-ct-rt]] === Compile-time and runtime properties @@ -279,12 +265,9 @@ These are compile-time properties that are applicable to a free function kernel, so they fit the model described above and may appear as the mandatory first argument of the [code]#SYCL_KHR_KERNEL# decoration. -As with all SYCL properties, properties can only be defined by the SYCL -implementation, so the SYCL specification (or an extension specification) -provides a formal specification for each property. -Apart from the _kind_ properties, any property that is applicable to a free -function kernel is defined by the SYCL implementation or by another extension -specification; this extension does not define any such property. +Apart from the _kind_ properties, this extension does not define any property +applicable to a free function kernel; any such property is defined by the SYCL +implementation or by another extension. [[sec:khr-free-function-kernels-arg-types]] == Restrictions on kernel argument types @@ -299,29 +282,18 @@ specification; this extension does not define any such property. A free function kernel passes each of its kernel arguments as a parameter of the function, rather than as a member of a function object or a lambda capture. -Each parameter of a free function kernel must have a type that is -<>. -An application can query whether a type [code]#T# is device copyable with the -core [code]#sycl::is_device_copyable_v# type trait. - -A free function kernel admits only device-copyable parameters. -The rules for parameter passing to kernels (<>) -permit two categories of kernel parameter type: any <> type, -and a separate set of special SYCL types — such as [code]#accessor#, -[code]#local_accessor#, the image accessors, [code]#stream#, [code]#reducer#, -and [code]#kernel_handler# — that are legal only because the implementation -passes them in a special, non-positional way. -A free function kernel does not provide this second allowance: only the first -category applies. -A program that uses any of these special types, or any other type that is not -<>, as a free function kernel parameter is ill formed. - -[_Note:_ A free function kernel receives every parameter positionally, as an -ordinary device-copyable kernel argument. -The special types above are not device copyable; they are legal parameters for -general SYCL kernels only through the separate special-passing allowance of -<>, which this extension does not provide. -_{endnote}_] +Each parameter must have a type that is <>; an application can +query this with the core [code]#sycl::is_device_copyable_v# type trait. + +The general rules for parameter passing to kernels +(<>) also permit a set of special SYCL types --- +such as [code]#accessor#, [code]#local_accessor#, the image accessors, +[code]#stream#, [code]#reducer#, and [code]#kernel_handler# --- which are legal +only because the implementation passes them in a special, non-positional way. +A free function kernel receives every parameter positionally and does not provide +this second allowance: a program that uses any of these special types, or any +other type that is not <>, as a free function kernel parameter +is ill formed. [[sec:khr-free-function-kernels-traits]] == Traits for kernel functions @@ -341,14 +313,12 @@ _{endnote}_] This extension defines traits that report, at compile time, whether a given [code]#Func# is a free function kernel (see <>) and, if so, what kind it has. -In each trait, [code]#Func# is the address of a function; the trait reports -[code]#true# only when [code]#Func# is the address of a function that is -decorated as a free function kernel of the relevant kind. - -These traits are usable in constant expressions and in SFINAE contexts at -compile time, so an application (or a software layer built on top of this -extension) can branch on, or constrain a template against, the kind of a -free function kernel without launching it. +In each trait, [code]#Func# is the address of a function, and the trait reports +[code]#true# only when [code]#Func# is decorated as a free function kernel of the +relevant kind. +These traits are usable in constant expressions and in SFINAE contexts, so an +application can branch on, or constrain a template against, the kind of a free +function kernel without launching it. {note}The kernel whose address is [code]#Func# need not be defined in the same translation unit as the use of the trait; the declaration decorated with @@ -444,12 +414,10 @@ constant expression. A free function kernel is launched by passing its address to one of the launch functions defined in this section. -The kernel's address is supplied through the [code]#kernel_function# handle, -which carries the address as a template parameter. -A launch function deduces the kernel function [code]#Func# from this handle, so -the kernel is identified by the _value_ [code]#sycl::khr::kernel_function# -that is passed as an argument --- not by an explicit template argument on the -launch function itself. +The address is supplied through the [code]#kernel_function# handle, which carries +it as a template parameter, so the kernel is identified by the _value_ +[code]#sycl::khr::kernel_function# passed as an argument rather than by an +explicit template argument on the launch function. .[apidef]#kernel_function# [source,role=synopsis,id=api:khr-free-function-kernels-kernel_function] @@ -534,8 +502,7 @@ kernel, using the number of work-items specified by the [code]#nd_range# Each value in the [code]#args# pack is passed to the corresponding parameter of [code]#Func#, converting it to the parameter's type if necessary. -The launch functions enqueue the kernel by its address; the kernel is selected -by the [code]#kernel_function# value, as in: +For example: [source] ---- @@ -546,14 +513,12 @@ The [code]#Constraints# above check, at compile time, only the kernel's _kind_ and _dimensionality_ and that the supplied arguments can be passed to [code]#Func#. They do not check the values carried by the kernel's compile-time properties, -such as a required work-group size: a required work-group size is compared -against the local size of the [code]#nd_range# at launch, and the local size of -an [code]#nd_range# is generally a run-time value, so this comparison is -performed at run time rather than at compile time. - -Any launch requirements that a free function kernel carries through its -compile-time properties are enforced at launch in the same way as for any other -SYCL kernel, as described in <>. +such as a required work-group size, because such a value is compared against the +local size of the [code]#nd_range#, which is generally a run-time value. + +Launch requirements that a free function kernel carries through its compile-time +properties are enforced at launch in the same way as for any other SYCL kernel +(<>). In particular, launching a kernel that declares a required work-group size with an [code]#nd_range# whose local size does not match that requirement throws a synchronous [code]#exception# with the [code]#errc::nd_range# error code, @@ -583,10 +548,8 @@ that can be passed through them.{endnote} The launch functions in <> have additional overloads that accept a _launch property list_ --- a <> properties list supplied at the launch -site. -The property list is passed as a separate argument, immediately before the -[code]#kernel_function# handle; it is not wrapped in any launch-configuration -type. +site, passed as a separate argument immediately before the +[code]#kernel_function# handle and not wrapped in any launch-configuration type. The properties in the list carry _runtime_ information that applies to the individual launch, as distinct from the compile-time properties that decorate the kernel itself (see <>). @@ -601,9 +564,7 @@ shape that future runtime launch properties will use. {note}A runtime launch property is one whose value is not known until the kernel is launched, so it cannot be carried by the compile-time decoration on the kernel. A likely future example, defined by another extension, is a dynamic work-group -local memory size, whose byte count is a run-time value supplied at launch. -This extension reserves the launch-property overloads so that such a property can -be added without changing the shape of the launch API.{endnote} +local memory size, whose byte count is a run-time value supplied at launch.{endnote} [[sec:khr-free-function-kernels-launch-properties-single_task]] === single_task From f5e2cd75701a2654f9312384b45e70cfea355ced Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 25 Jun 2026 16:43:42 -0700 Subject: [PATCH 10/16] Cleanup --- .../sycl_khr_free_function_kernels.adoc | 170 +++++++----------- 1 file changed, 69 insertions(+), 101 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 8158fdf27..da46e732b 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -2,13 +2,14 @@ = sycl_khr_free_function_kernels This extension introduces _free function kernels_. -A free function kernel is an ordinary C++ function, defined at namespace scope or -as a static member function, that is decorated with the [code]#SYCL_KHR_KERNEL# +A free function kernel is an ordinary C++ function, defined at namespace scope, +that is decorated with the [code]#SYCL_KHR_KERNEL# macro so that it becomes a device kernel entry point. Its kernel arguments are the function's parameters, rather than the captures of a lambda or the members of a function object, which gives them a defined order. -Because it is a named function, the kernel is identified by its address, so it can -be launched and queried by reference to the function rather than only as a lambda. +Because it is a named function, the kernel is identified by the function itself, +so it can be launched and queried by reference to the function rather than only as +a lambda. [[sec:khr-free-function-kernels-dependencies]] == Dependencies @@ -26,10 +27,12 @@ infrastructure it uses to associate properties with a free function kernel (see <>). This extension depends on the -<> extension, which -provides the in-kernel queries (such as [code]#sycl::khr::this_nd_item#) that the -body of a [code]#sycl::khr::nd_kernel# free function kernel uses to obtain -its work-item's position in the iteration space. +<> extension. A free +function kernel obtains its work-item's position in the iteration space, when it +needs one, through the in-kernel queries provided by that extension (such as +[code]#sycl::khr::this_nd_item#). A kernel that does not need a position --- a +[code]#sycl::khr::single_task_kernel#, or a [code]#sycl::khr::nd_kernel# +whose body never queries one --- uses no such query. {note}At the time of writing, [code]#sycl_khr_properties# is a proposed extension (SYCL specification PR #980) and is not yet part of the SYCL specification. @@ -63,8 +66,7 @@ below. // Ported from experimental "=== Defining a free function kernel" (:156-205, // :324-330). Reframed onto the mandatory-kind macro SYCL_KHR_KERNEL(kind, ...) // (project memory step 12) rather than the experimental -// SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((property)) decoration. Arg-type rule is NOT -// restated here — cross-refs <>. +// SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((property)) decoration. // SYCL_EXTERNAL is left as an OPEN note (not a normative rule, per project // memory). Call-as-a-function is stated as UB only (impl is not claimed to // reject it). See playground/ffk-khr-coverage.md. @@ -109,16 +111,15 @@ A program that violates any of them is ill formed unless stated otherwise. Invoking the macro with no kind --- [code]#SYCL_KHR_KERNEL()# --- is ill formed. -* The function must be declared at namespace scope, or at class scope as a - static member function. +* The function must be declared at namespace scope. * The function's return type must be [code]#void#. * The function must not accept a variadic argument list. -* Each of the function's parameters must have a type that is permitted as a free - function kernel argument, as defined in - <>. +* Each of the function's parameters must have a type that is <>; + an application can query this with the core + [code]#sycl::is_device_copyable_v# type trait. * No declaration of the function may specify a default argument for any parameter. @@ -139,30 +140,21 @@ A program that violates any of them is ill formed unless stated otherwise. unit that declares the same function must decorate it identically (the same kind and the same property arguments). -* The way the function's body obtains its iteration ID must be consistent with - the declared kind: a function whose kind is - [code]#sycl::khr::nd_kernel# must obtain its iteration ID for an - [code]#nd_range# of [code]#Dims# dimensions, and a function whose kind is - [code]#sycl::khr::single_task_kernel# must not obtain an iteration ID. - A program that violates this rule is ill formed, no diagnostic required. - -[_Note:_ An implementation may, but need not, diagnose an inconsistency between -the function's body and its declared kind. -Such a diagnosis is not always possible, for example when the body is provided in -another translation unit or as a precompiled device object. -_{endnote}_] +The way the function's body obtains its position in the iteration space must be +consistent with the declared kind: a function whose kind is +[code]#sycl::khr::nd_kernel# obtains its position for an +[code]#nd_range# of [code]#Dims# dimensions, while a function whose kind is +[code]#sycl::khr::single_task_kernel# executes as a single work-item, with no +iteration space, and must not query an iteration position. +A program that violates this rule is ill formed, no diagnostic required. A function decorated with [code]#SYCL_KHR_KERNEL# is a kernel entry point. Calling it as an ordinary function, in either host or device code, results in undefined behavior. -The function's address may still be taken; this address identifies the kernel and -is the value passed to the launch and query APIs described below (the -[code]#Func# parameter). -[_Note:_ Because such a call is undefined behavior rather than ill formed, an -implementation is not required to diagnose it, though it is permitted to diagnose -one that appears in host code as a quality-of-implementation matter. -_{endnote}_] +The function itself identifies the kernel: the launch and query APIs described +below are parameterized on it through a non-type template parameter +[code]#Func# (for example, [code]#sycl::khr::kernel_function#). // OPEN (raise in PR discussion, not a normative rule): may a free function // kernel entry point be declared SYCL_EXTERNAL? Tension — the host needs the @@ -203,16 +195,10 @@ kernel: [[sec:khr-free-function-kernels-properties-tag]] === The free_function_kernel property tag -The [code]#sycl_khr_properties# extension expresses the set of classes that a -property may be used with through the [api]#khr::is_property_for# trait, which is -keyed on a _class_ type: [code]#khr::is_property_for_v# is [code]#true# -when the property [code]#P# can be used with the class [code]#Class#. -A free function kernel, however, is a decorated _function_ rather than an object -of some class, so there is no class type to use as the [code]#Class# argument. -This extension therefore introduces an empty tag type that stands in for the -free function kernel as that class, so the existing -<> machinery is reused unchanged with no -new applicability trait. +The tag type [code]#sycl::khr::free_function_kernel# identifies a free function +kernel for the purpose of property applicability: a property [code]#P# is +applicable to a free function kernel if +[code]#khr::is_property_for_v# is [code]#true#. ''' @@ -229,11 +215,7 @@ struct free_function_kernel; _Remarks:_ [code]#free_function_kernel# is an incomplete tag type that is never instantiated. It serves only as the [code]#Class# argument of the -<> property traits, identifying the -properties that may be associated with a free function kernel. - -A property [code]#P# is _applicable to a free function kernel_ if, and only if, -[code]#khr::is_property_for_v# is [code]#true#. +<> property traits. [[sec:khr-free-function-kernels-properties-ct-rt]] === Compile-time and runtime properties @@ -258,43 +240,17 @@ compile-time property or a runtime property, as classified by the [[sec:khr-free-function-kernels-properties-kind]] === Properties defined by this extension -This extension defines only the _kind_ properties, +This extension defines the compile-time _kind_ properties, [code]#sycl::khr::single_task_kernel# and [code]#sycl::khr::nd_kernel# (see <> for their meaning). -These are compile-time properties that are applicable to a free function kernel, -so they fit the model described above and may appear as the mandatory first +These are applicable to a free function kernel, +and appear as the mandatory first argument of the [code]#SYCL_KHR_KERNEL# decoration. Apart from the _kind_ properties, this extension does not define any property applicable to a free function kernel; any such property is defined by the SYCL implementation or by another extension. -[[sec:khr-free-function-kernels-arg-types]] -== Restrictions on kernel argument types - -// COVERAGE: arg type rule = device-copyable, the SAME rule as any SYCL kernel -// argument (<>); reuse core is_device_copyable_v -// rather than a new trait (project memory "ARG-TYPE TWO-TIER DECISION": spec -// tier = DEVICE COPYABLE; the tighter cross-ABI ceiling is off-spec flag-mode -// dpcpp, NOT in this spec). Prototype steps 2/3 (is_valid_kernel_arg_v + -// narrowing) implement the off-spec flag tier, not this section. See -// playground/ffk-khr-coverage.md and ffk-khr-pr-step-C-args-RESULT.md. - -A free function kernel passes each of its kernel arguments as a parameter of the -function, rather than as a member of a function object or a lambda capture. -Each parameter must have a type that is <>; an application can -query this with the core [code]#sycl::is_device_copyable_v# type trait. - -The general rules for parameter passing to kernels -(<>) also permit a set of special SYCL types --- -such as [code]#accessor#, [code]#local_accessor#, the image accessors, -[code]#stream#, [code]#reducer#, and [code]#kernel_handler# --- which are legal -only because the implementation passes them in a special, non-positional way. -A free function kernel receives every parameter positionally and does not provide -this second allowance: a program that uses any of these special types, or any -other type that is not <>, as a free function kernel parameter -is ill formed. - [[sec:khr-free-function-kernels-traits]] == Traits for kernel functions @@ -561,11 +517,6 @@ This version of the extension defines no such property: the in-scope set of runtime launch properties is empty, and the overloads below exist to reserve the shape that future runtime launch properties will use. -{note}A runtime launch property is one whose value is not known until the kernel -is launched, so it cannot be carried by the compile-time decoration on the kernel. -A likely future example, defined by another extension, is a dynamic work-group -local memory size, whose byte count is a run-time value supplied at launch.{endnote} - [[sec:khr-free-function-kernels-launch-properties-single_task]] === single_task @@ -646,12 +597,6 @@ _Effects:_ Equivalent to the [code]#nd_launch# overload without a property list (see <>), additionally applying the runtime properties in [code]#props# to the launch. -{note}Because this version of the extension defines no runtime property that is -applicable to a free function kernel, the only property list that satisfies the -[code]#Constraints# above is the empty properties list. -The overloads are nonetheless specified so that the launch API does not change -shape when a future extension adds a runtime launch property.{endnote} - [[sec:khr-free-function-kernels-example]] == Example @@ -699,9 +644,12 @@ int main() { } ---- -A free function kernel may also be a function template. -The address passed through [code]#kernel_function# is then the address of a -particular specialization. +A free function kernel may also be a function template, or one of several +overloads. +Because the kernel is identified by the non-type template parameter +[code]#Func#, the program must name the exact function whose address is taken: +for a template, pass the address of a specific specialization; for an overload +set, cast to the desired function type to select the overload. [source,,linenums] ---- @@ -710,27 +658,47 @@ particular specialization. constexpr size_t N = 1024; constexpr size_t WGSIZE = 32; +// A templated free function kernel. template SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) -void fill(T value, T *data) { +void iota(T start, T *p) { + size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); + p[i] = start + static_cast(i); +} + +// Two overloads of a free function kernel. +SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +void ping(float *p) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); - data[i] = value; + p[i] = 1.0f; +} + +SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +void ping(int *p) { + size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); + p[i] = 1; } int main() { sycl::queue q; - int *data = sycl::malloc_shared(N, q); + float *fptr = sycl::malloc_shared(N, q); + int *iptr = sycl::malloc_shared(N, q); + sycl::nd_range<1> ndr{N, WGSIZE}; // For a templated kernel, pass the address of a specific instantiation. - sycl::khr::nd_launch(q, sycl::nd_range<1>{N, WGSIZE}, - sycl::khr::kernel_function>, 42, data); - q.wait(); + sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function>, 3.14f, fptr); + sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function>, 3, iptr); - for (size_t i = 0; i < N; ++i) - assert(data[i] == 42); + // For an overloaded kernel, cast to the desired function type to select the + // overload. + sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function<(void(*)(float*))ping>, fptr); + sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function<(void(*)(int*))ping>, iptr); - sycl::free(data, q); + q.wait(); + + sycl::free(fptr, q); + sycl::free(iptr, q); return 0; } ---- From 84ba00c46beaa2e590cc7c26400da7a6757739df Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Tue, 7 Jul 2026 13:20:20 -0700 Subject: [PATCH 11/16] Address Greg's comments --- .../sycl_khr_free_function_kernels.adoc | 458 ++++-------------- 1 file changed, 97 insertions(+), 361 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index da46e732b..5140c5f25 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -3,8 +3,9 @@ This extension introduces _free function kernels_. A free function kernel is an ordinary C++ function, defined at namespace scope, -that is decorated with the [code]#SYCL_KHR_KERNEL# -macro so that it becomes a device kernel entry point. +that is decorated with one of the [code]#SYCL_KHR_ND_KERNEL# or +[code]#SYCL_KHR_TASK_KERNEL# macros so that it becomes a device kernel entry +point. Its kernel arguments are the function's parameters, rather than the captures of a lambda or the members of a function object, which gives them a defined order. Because it is a named function, the kernel is identified by the function itself, @@ -14,33 +15,13 @@ a lambda. [[sec:khr-free-function-kernels-dependencies]] == Dependencies -// COVERAGE: this section OWNS the #980 (sycl_khr_properties) unmerged -// forward-dependency caveat for the whole document — the editorial note below is -// the single place that explains why every <> cross-ref -// in this extension dangles in a build until #980 lands. Do not restate the -// property model here (see <>). See -// playground/ffk-khr-pr-step-C7-dependencies.md. - -This extension depends on the -<> extension, whose property -infrastructure it uses to associate properties with a free function kernel (see -<>). - This extension depends on the <> extension. A free function kernel obtains its work-item's position in the iteration space, when it needs one, through the in-kernel queries provided by that extension (such as [code]#sycl::khr::this_nd_item#). A kernel that does not need a position --- a -[code]#sycl::khr::single_task_kernel#, or a [code]#sycl::khr::nd_kernel# -whose body never queries one --- uses no such query. - -{note}At the time of writing, [code]#sycl_khr_properties# is a proposed extension -(SYCL specification PR #980) and is not yet part of the SYCL specification. -Until it is incorporated, the cross-references to it in this extension do not -resolve.{endnote} - -Some features of this extension are only available when a SYCL implementation -conforms to {cpp17} or later. +single-task free function kernel, or an ND-range free function kernel whose body +never queries one --- uses no such query. [[sec:khr-free-function-kernels-feature-test]] == Feature test macro @@ -61,35 +42,28 @@ below. [[sec:khr-free-function-kernels-defining]] == Defining a free function kernel -// COVERAGE: step 1 (variadic SYCL_KHR_KERNEL macro, khr_free_kernel_macro.cpp) + -// step 12 (mandatory kind / empty-macro-is-error, khr_free_kernel_macro.cpp). -// Ported from experimental "=== Defining a free function kernel" (:156-205, -// :324-330). Reframed onto the mandatory-kind macro SYCL_KHR_KERNEL(kind, ...) -// (project memory step 12) rather than the experimental -// SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((property)) decoration. -// SYCL_EXTERNAL is left as an OPEN note (not a normative rule, per project -// memory). Call-as-a-function is stated as UB only (impl is not claimed to -// reject it). See playground/ffk-khr-coverage.md. - A free function kernel is an ordinary C++ function whose declaration is decorated -with the [code]#SYCL_KHR_KERNEL# macro. -The first argument to the macro is the kernel's _kind_, which is mandatory. -The kind is one of the following: +with one of the [code]#SYCL_KHR_ND_KERNEL# or [code]#SYCL_KHR_TASK_KERNEL# +macros. +The kernel's _kind_ is determined by which macro is used: -* [code]#sycl::khr::single_task_kernel# --- the function is launched as a single - task, without any iteration space; or +* [code]#SYCL_KHR_ND_KERNEL(N)# --- the function is launched over an + [code]#nd_range# iteration space of [code]#N# dimensions; or -* [code]#sycl::khr::nd_kernel# --- the function is launched over an - [code]#nd_range# iteration space of [code]#Dims# dimensions. +* [code]#SYCL_KHR_TASK_KERNEL()# --- the function is launched as a single task, + without any iteration space. -Any remaining arguments to the macro are optional compile-time properties that -are applicable to a free function kernel, as described in -<>. +Neither macro takes any arguments beyond those shown above, though a future +extension may define more. [source,role=synopsis,id=api:khr-free-function-kernels-macro] ---- -// Decorate a function declaration to define a free function kernel. -#define SYCL_KHR_KERNEL(kind, ...) /* see below */ +// Decorate a function declaration to define an ND-range free function kernel +// over N dimensions. +#define SYCL_KHR_ND_KERNEL(N) /* see below */ + +// Decorate a function declaration to define a single-task free function kernel. +#define SYCL_KHR_TASK_KERNEL() /* see below */ ---- The first declaration of the function in the translation unit must be the @@ -98,7 +72,7 @@ For example: [source] ---- -SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +SYCL_KHR_ND_KERNEL(1) void scale(float factor, float *data) { // ... } @@ -107,9 +81,7 @@ void scale(float factor, float *data) { A free function kernel must obey all of the following rules. A program that violates any of them is ill formed unless stated otherwise. -* The macro's kind argument must be present. - Invoking the macro with no kind --- [code]#SYCL_KHR_KERNEL()# --- is ill - formed. +* The function must be decorated with exactly one of the two macros. * The function must be declared at namespace scope. @@ -127,28 +99,31 @@ A program that violates any of them is ill formed unless stated otherwise. * The decoration must appear on the first declaration of the function in the translation unit. A redeclaration of the function may also be decorated, provided it is decorated - with the same kind (and, where present, the same property arguments). + with the same macro (the same kind, and for [code]#SYCL_KHR_ND_KERNEL# the same + [code]#N#). The effect is the same whether or not a redeclaration is decorated. -* The same function must be decorated with exactly one kind. - Decorating the same function with more than one kind, or with the same kind but - conflicting property arguments, is ill formed. - Decorating it more than once with the same kind and the same property arguments - is permitted and has the same effect as decorating it once. - * If the function is decorated in one translation unit, every other translation - unit that declares the same function must decorate it identically (the same - kind and the same property arguments). + unit that declares the same function must decorate it with the same macro (the + same kind, and for [code]#SYCL_KHR_ND_KERNEL# the same [code]#N#). + +{note}Some SYCL types that are legal parameters for an ordinary SYCL kernel are +not legal parameters for a free function kernel, because a free function kernel +requires each parameter to be device-copyable (see the rule above). +For example, [code]#accessor#, [code]#local_accessor#, image accessors, +[code]#stream#, and [code]#reducer# are legal kernel parameters under +<> but are not device-copyable and so cannot be +used as free function kernel parameters.{endnote} The way the function's body obtains its position in the iteration space must be -consistent with the declared kind: a function whose kind is -[code]#sycl::khr::nd_kernel# obtains its position for an -[code]#nd_range# of [code]#Dims# dimensions, while a function whose kind is -[code]#sycl::khr::single_task_kernel# executes as a single work-item, with no -iteration space, and must not query an iteration position. -A program that violates this rule is ill formed, no diagnostic required. - -A function decorated with [code]#SYCL_KHR_KERNEL# is a kernel entry point. +consistent with the kind implied by the macro used: a function decorated with +[code]#SYCL_KHR_ND_KERNEL(N)# obtains its position for an [code]#nd_range# of +[code]#N# dimensions, while a function decorated with +[code]#SYCL_KHR_TASK_KERNEL()# executes as a single work-item, with no iteration +space, and must not query an iteration position. +A program that violates this rule is ill formed, no diagnostic is required. + +A function decorated with one of these macros is a kernel entry point. Calling it as an ordinary function, in either host or device code, results in undefined behavior. @@ -156,116 +131,9 @@ The function itself identifies the kernel: the launch and query APIs described below are parameterized on it through a non-type template parameter [code]#Func# (for example, [code]#sycl::khr::kernel_function#). -// OPEN (raise in PR discussion, not a normative rule): may a free function -// kernel entry point be declared SYCL_EXTERNAL? Tension — the host needs the -// symbol visible (to name/launch it) but never calls it, while SYCL_EXTERNAL is -// a device cross-TU linkage contract that is meaningless for an entry point -// (calling a kernel from device code is UB). The experimental spec (:204) -// contemplates SYCL_EXTERNAL free function kernel bodies; whether the KHR should -// permit, forbid, or stay silent on this is deferred to the PR. Do NOT add a -// normative SYCL_EXTERNAL rule here until that is resolved. - -[[sec:khr-free-function-kernels-properties]] -== Kernel properties - -// COVERAGE: the _kind_ properties (single_task_kernel / nd_kernel) are -// PROVEN (step 1 macro + step 12 mandatory-kind validator, -// khr_free_kernel_macro.cpp). The free_function_kernel TAG TYPE and the -// is_property_for-keyed applicability model are a NEW SPEC CONSTRUCT (design -// proposal, project memory "FFK NEEDS A PROPERTY-TAG TYPE"): the prototype used -// khr::detail introspection of the property bundle, NOT a public tag. The -// compile-time vs runtime split reuses Greg's is_property_key_compile_time -// discriminator (step 6, khr_launch_properties*.cpp). DEPENDS on the unmerged -// sycl_khr_properties (#980) — the is_property_for / is_property_key_compile_time -// cross-refs below dangle (sec:khr-properties) until it lands. See -// playground/ffk-khr-coverage.md and ffk-khr-pr-step-C5-property-model.md. - -A free function kernel can carry _properties_, using the property infrastructure -defined by <>. -There are two contexts in which a property is associated with a free function -kernel: - -* _Compile-time properties_ are placed on the kernel by listing them as the - remaining arguments of the [code]#SYCL_KHR_KERNEL# decoration (see - <>). - -* _Runtime properties_ are supplied when the kernel is launched, through the - launch property list (see <>). - -[[sec:khr-free-function-kernels-properties-tag]] -=== The free_function_kernel property tag - -The tag type [code]#sycl::khr::free_function_kernel# identifies a free function -kernel for the purpose of property applicability: a property [code]#P# is -applicable to a free function kernel if -[code]#khr::is_property_for_v# is [code]#true#. - -''' - -.[apidef]#khr::free_function_kernel# -[source,role=synopsis,id=api:khr-free-function-kernels-tag] ----- -namespace sycl::khr { - -struct free_function_kernel; - -} // namespace sycl::khr ----- - -_Remarks:_ [code]#free_function_kernel# is an incomplete tag type that is never -instantiated. -It serves only as the [code]#Class# argument of the -<> property traits. - -[[sec:khr-free-function-kernels-properties-ct-rt]] -=== Compile-time and runtime properties - -Whether a property is associated with a free function kernel through the -decoration or through the launch property list is determined by whether it is a -compile-time property or a runtime property, as classified by the -[api]#khr::is_property_key_compile_time# trait of -<>. - -* The [code]#SYCL_KHR_KERNEL# decoration accepts a property [code]#P# only if - [code]#P# is applicable to a free function kernel and the key of [code]#P# is a - compile-time property, that is - [code]#khr::is_property_key_compile_time_v# is [code]#true#. - -* The launch property list (see - <>) accepts a property - [code]#P# only if [code]#P# is applicable to a free function kernel and the key - of [code]#P# is a runtime property, that is - [code]#khr::is_property_key_compile_time_v# is [code]#false#. - -[[sec:khr-free-function-kernels-properties-kind]] -=== Properties defined by this extension - -This extension defines the compile-time _kind_ properties, -[code]#sycl::khr::single_task_kernel# and [code]#sycl::khr::nd_kernel# -(see <> for their meaning). -These are applicable to a free function kernel, -and appear as the mandatory first -argument of the [code]#SYCL_KHR_KERNEL# decoration. - -Apart from the _kind_ properties, this extension does not define any property -applicable to a free function kernel; any such property is defined by the SYCL -implementation or by another extension. - [[sec:khr-free-function-kernels-traits]] == Traits for kernel functions -// COVERAGE: step 4 (is_kernel / is_nd_kernel / is_single_task_kernel traits -// queryable at host compile time off the decorated declaration, NO integration -// header — sycl-property-builtins.cpp, khr_kernel_property_traits.cpp). Ported -// from experimental "=== New traits for kernel functions" (:336-417); renamed -// is_nd_range_kernel -> is_nd_kernel (khr kind name) and moved to sycl::khr. -// Spec states ONLY the OBSERVABLE contract ("usable in constant expressions / -// SFINAE at compile time"); the builtin / integration-header / device-compile -// mechanism is an implementation detail and is deliberately not specified. -// A general has_property / get_property query surface is DEFERRED (it belongs -// with sycl_khr_properties, #980) — only the three kind/dimensionality traits -// are exposed here, matching the experimental spec. See coverage map. - This extension defines traits that report, at compile time, whether a given [code]#Func# is a free function kernel (see <>) and, if so, what kind it has. @@ -277,8 +145,8 @@ application can branch on, or constrain a template against, the kind of a free function kernel without launching it. {note}The kernel whose address is [code]#Func# need not be defined in the same -translation unit as the use of the trait; the declaration decorated with -[code]#SYCL_KHR_KERNEL# is sufficient. +translation unit as the use of the trait; the decorated declaration (with +[code]#SYCL_KHR_ND_KERNEL# or [code]#SYCL_KHR_TASK_KERNEL#) is sufficient. How an implementation makes the decoration observable in a constant expression is unspecified.{endnote} @@ -322,52 +190,39 @@ inline constexpr bool is_nd_kernel_v = is_nd_kernel::value; ---- _Returns:_ [code]#is_nd_kernel::value# is [code]#true# if -[code]#Func# is the address of a function whose kind is -[code]#sycl::khr::nd_kernel#, and [code]#false# otherwise. +[code]#Func# is the address of a function that is decorated as an ND-range free +function kernel over [code]#Dims# dimensions (with +[code]#SYCL_KHR_ND_KERNEL(Dims)#), and [code]#false# otherwise. _Remarks:_ The expression [code]#is_nd_kernel_v# is usable in a constant expression. ''' -.[apidef]#is_single_task_kernel# -[source,role=synopsis,id=api:khr-free-function-kernels-is_single_task_kernel] +.[apidef]#is_task_kernel# +[source,role=synopsis,id=api:khr-free-function-kernels-is_task_kernel] ---- namespace sycl::khr { template -struct is_single_task_kernel; +struct is_task_kernel; template -inline constexpr bool is_single_task_kernel_v = is_single_task_kernel::value; +inline constexpr bool is_task_kernel_v = is_task_kernel::value; } // namespace sycl::khr ---- -_Returns:_ [code]#is_single_task_kernel::value# is [code]#true# if -[code]#Func# is the address of a function whose kind is -[code]#sycl::khr::single_task_kernel#, and [code]#false# otherwise. +_Returns:_ [code]#is_task_kernel::value# is [code]#true# if +[code]#Func# is the address of a function that is decorated as a single-task free +function kernel (with [code]#SYCL_KHR_TASK_KERNEL#), and [code]#false# otherwise. -_Remarks:_ The expression [code]#is_single_task_kernel_v# is usable in a +_Remarks:_ The expression [code]#is_task_kernel_v# is usable in a constant expression. [[sec:khr-free-function-kernels-launch]] == Launching a free function kernel -// PORT: experimental "=== New free functions to launch a kernel" (:418-502) + -// "=== Enqueuing a free function kernel and setting parameter values" -// (:800-841). Define khr::nd_launch / khr::single_task (queue + handler forms), -// kernel identity = khr::kernel_function value, both return void. -// Compile-time checks: kind/dimensionality (is_nd_range_kernel_v) + -// arg-type (is_valid_kernel_arg_v) via Constraints. NO compile-time -// work-group-size value check (Greg feedback — constant nd_ranges are rare). -// Runtime must-match (work_group_size etc.) noted as runtime exceptions. -// COVERAGE: step 5 (checked nd_launch/single_task; khr_checked_launch.cpp, -// khr_checked_launch_errors.cpp) + step 7/8 (runtime must-match throws -// errc::nd_range, decoration propagated onto wrapper; e2e khr_must_match.cpp, -// khr_must_match_negative.cpp on PVC — work_group_size proven, max_*/sub_group -// designed-not-e2e-tested). See coverage map. - A free function kernel is launched by passing its address to one of the launch functions defined in this section. The address is supplied through the [code]#kernel_function# handle, which carries @@ -392,32 +247,32 @@ inline constexpr kernel_function_s kernel_function; _Remarks:_ [code]#Func# is the address of a function that is decorated as a free function kernel (see <>). The variable template [code]#kernel_function# is the value passed to -[code]#single_task# and [code]#nd_launch# to identify the kernel to launch. +[code]#launch_task# and [code]#launch_grouped# to identify the kernel to launch. {note}For a free function kernel that is a template, the address of a particular specialization is used, for example [code]#kernel_function>#.{endnote} -[[sec:khr-free-function-kernels-launch-single_task]] -=== single_task +[[sec:khr-free-function-kernels-launch-launch_task]] +=== launch_task ''' -.[apidef]#single_task# -[source,role=synopsis,id=api:khr-free-function-kernels-single_task] +.[apidef]#launch_task# +[source,role=synopsis,id=api:khr-free-function-kernels-launch_task] ---- namespace sycl::khr { template -void single_task(queue q, kernel_function_s k, Args&&... args); +void launch_task(const queue &q, kernel_function_s k, Args&&... args); template -void single_task(handler &h, kernel_function_s k, Args&&... args); +void launch_task(handler &h, kernel_function_s k, Args&&... args); } // namespace sycl::khr ---- -_Constraints:_ Available only if [code]#is_single_task_kernel_v# is +_Constraints:_ Available only if [code]#is_task_kernel_v# is [code]#true# and [code]#std::is_invocable_v# is [code]#true#. @@ -426,23 +281,23 @@ _Effects:_ Enqueues the free function kernel [code]#Func# to the queue Each value in the [code]#args# pack is passed to the corresponding parameter of [code]#Func#, converting it to the parameter's type if necessary. -[[sec:khr-free-function-kernels-launch-nd_launch]] -=== nd_launch +[[sec:khr-free-function-kernels-launch-launch_grouped]] +=== launch_grouped ''' -.[apidef]#nd_launch# -[source,role=synopsis,id=api:khr-free-function-kernels-nd_launch] +.[apidef]#launch_grouped# +[source,role=synopsis,id=api:khr-free-function-kernels-launch_grouped] ---- namespace sycl::khr { template -void nd_launch(queue q, nd_range r, - kernel_function_s k, Args&&... args); +void launch_grouped(const queue &q, range global, range local, + kernel_function_s k, Args&&... args); template -void nd_launch(handler &h, nd_range r, - kernel_function_s k, Args&&... args); +void launch_grouped(handler &h, range global, range local, + kernel_function_s k, Args&&... args); } // namespace sycl::khr ---- @@ -453,8 +308,8 @@ _Constraints:_ Available only if [code]#is_nd_kernel_v# is _Effects:_ Enqueues the free function kernel [code]#Func# to the queue [code]#q# (or to the command group of the handler [code]#h#) as an ND-range -kernel, using the number of work-items specified by the [code]#nd_range# -[code]#r#. +kernel, over the iteration space defined by the global range [code]#global# with +work-group (local) size [code]#local#. Each value in the [code]#args# pack is passed to the corresponding parameter of [code]#Func#, converting it to the parameter's type if necessary. @@ -462,149 +317,29 @@ For example: [source] ---- -sycl::khr::nd_launch(q, r, sycl::khr::kernel_function, factor, data); +sycl::khr::launch_grouped(q, global, local, sycl::khr::kernel_function, factor, data); ---- The [code]#Constraints# above check, at compile time, only the kernel's _kind_ and _dimensionality_ and that the supplied arguments can be passed to [code]#Func#. -They do not check the values carried by the kernel's compile-time properties, -such as a required work-group size, because such a value is compared against the -local size of the [code]#nd_range#, which is generally a run-time value. - -Launch requirements that a free function kernel carries through its compile-time -properties are enforced at launch in the same way as for any other SYCL kernel -(<>). -In particular, launching a kernel that declares a required work-group size with -an [code]#nd_range# whose local size does not match that requirement throws a -synchronous [code]#exception# with the [code]#errc::nd_range# error code, -exactly as for a SYCL kernel decorated with [code]#reqd_work_group_size#. - -{note}The launch functions above do not accept a launch property list. -Overloads that do are defined in -<>; this version of the -extension reserves the shape of those overloads but defines no runtime property -that can be passed through them.{endnote} - -[[sec:khr-free-function-kernels-launch-properties]] -== Launch properties - -// COVERAGE: step 6 (runtime-property launch overload, CT props rejected; -// khr_launch_properties.cpp, khr_launch_properties_errors.cpp). The overload -// shape is RESERVED on BOTH nd_launch and single_task (project memory -// "LAUNCH-PROPERTY SECTION DESIGN"): symmetric and future-proof, even though -// single_task has no in-scope launch property in this version. The CT/RT split -// reuses Greg's is_property_key_compile_time discriminator; applicability reuses -// the free_function_kernel tag (C5). v1 in-scope payload is EMPTY -// (work_group_scratch_size is deferred to another extension) — this version -// defines no runtime launch property. DEPENDS on the unmerged sycl_khr_properties -// (#980): is_property_list_for / is_property_key_compile_time / sec:khr-properties -// dangle until it lands. See playground/ffk-khr-coverage.md. - -The launch functions in <> have -additional overloads that accept a _launch property list_ --- a -<> properties list supplied at the launch -site, passed as a separate argument immediately before the -[code]#kernel_function# handle and not wrapped in any launch-configuration type. -The properties in the list carry _runtime_ information that applies to the -individual launch, as distinct from the compile-time properties that decorate the -kernel itself (see <>). - -A launch property list accepts only properties that are _applicable to a free -function kernel_ (see <>) and whose -keys are _runtime properties_. -This version of the extension defines no such property: the in-scope set of -runtime launch properties is empty, and the overloads below exist to reserve the -shape that future runtime launch properties will use. - -[[sec:khr-free-function-kernels-launch-properties-single_task]] -=== single_task +They do not check the required work-group size of the kernel, if any, because +that value is compared against the local range, which is generally a run-time +value. -''' - -.[apidef]#single_task# -[source,role=synopsis,id=api:khr-free-function-kernels-single_task-props] ----- -namespace sycl::khr { - -template -void single_task(queue q, Properties props, - kernel_function_s k, Args&&... args); - -template -void single_task(handler &h, Properties props, - kernel_function_s k, Args&&... args); - -} // namespace sycl::khr ----- - -_Constraints:_ Available only if all of the following hold: - -* [code]#Properties# is a <> properties - list, that is [code]#khr::is_property_list_for_v# - is [code]#true#. - This requires every property in [code]#props# to be applicable to a free - function kernel (see <>). -* For every property in [code]#props#, the key of that property is a runtime - property, that is [code]#khr::is_property_key_compile_time_v# is - [code]#false#. - Supplying a compile-time property in the launch property list is ill formed. -* [code]#is_single_task_kernel_v# is [code]#true# and - [code]#std::is_invocable_v# is [code]#true#, as for - the launch function without a property list. - -_Effects:_ Equivalent to the [code]#single_task# overload without a property list -(see <>), additionally applying -the runtime properties in [code]#props# to the launch. - -[[sec:khr-free-function-kernels-launch-properties-nd_launch]] -=== nd_launch - -''' - -.[apidef]#nd_launch# -[source,role=synopsis,id=api:khr-free-function-kernels-nd_launch-props] ----- -namespace sycl::khr { - -template -void nd_launch(queue q, nd_range r, Properties props, - kernel_function_s k, Args&&... args); - -template -void nd_launch(handler &h, nd_range r, Properties props, - kernel_function_s k, Args&&... args); - -} // namespace sycl::khr ----- - -_Constraints:_ Available only if all of the following hold: - -* [code]#Properties# is a <> properties - list, that is [code]#khr::is_property_list_for_v# - is [code]#true#. - This requires every property in [code]#props# to be applicable to a free - function kernel (see <>). -* For every property in [code]#props#, the key of that property is a runtime - property, that is [code]#khr::is_property_key_compile_time_v# is - [code]#false#. - Supplying a compile-time property in the launch property list is ill formed. -* [code]#is_nd_kernel_v# is [code]#true# and - [code]#std::is_invocable_v# is [code]#true#, as for - the launch function without a property list. - -_Effects:_ Equivalent to the [code]#nd_launch# overload without a property list -(see <>), additionally applying -the runtime properties in [code]#props# to the launch. +A free function kernel is subject to the same launch-time requirements as any +other SYCL kernel (<>). +In particular, launching a kernel that requires a specific work-group size with a +local range that does not match that requirement throws a synchronous +[code]#exception# with the [code]#errc::nd_range# error code, exactly as for a +SYCL kernel decorated with [code]#reqd_work_group_size#. [[sec:khr-free-function-kernels-example]] == Example -// COVERAGE: steps 5 + 7 (positive e2e, khr_must_match.cpp on PVC). See -// playground/ffk-khr-coverage.md. - The example below defines a free function kernel at namespace scope, decorates it -with [code]#SYCL_KHR_KERNEL#, and launches it with [code]#sycl::khr::nd_launch#. +with [code]#SYCL_KHR_ND_KERNEL#, and launches it with +[code]#sycl::khr::launch_grouped#. The kernel obtains its work-item position with the <> query [code]#sycl::khr::this_nd_item#, rather than receiving an [code]#nd_item# @@ -612,13 +347,14 @@ parameter. [source,,linenums] ---- +#include #include constexpr size_t N = 1024; constexpr size_t WGSIZE = 32; // A free function kernel: an ordinary function, decorated with the kind. -SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +SYCL_KHR_ND_KERNEL(1) void scale(float factor, float *data) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); data[i] *= factor; @@ -632,8 +368,8 @@ int main() { data[i] = static_cast(i); // Identify the kernel by the function itself and launch it. - sycl::khr::nd_launch(q, sycl::nd_range<1>{N, WGSIZE}, - sycl::khr::kernel_function, 2.0f, data); + sycl::khr::launch_grouped(q, sycl::range<1>{N}, sycl::range<1>{WGSIZE}, + sycl::khr::kernel_function, 2.0f, data); q.wait(); for (size_t i = 0; i < N; ++i) @@ -660,20 +396,20 @@ constexpr size_t WGSIZE = 32; // A templated free function kernel. template -SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +SYCL_KHR_ND_KERNEL(1) void iota(T start, T *p) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); p[i] = start + static_cast(i); } // Two overloads of a free function kernel. -SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +SYCL_KHR_ND_KERNEL(1) void ping(float *p) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); p[i] = 1.0f; } -SYCL_KHR_KERNEL(sycl::khr::nd_kernel<1>) +SYCL_KHR_ND_KERNEL(1) void ping(int *p) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); p[i] = 1; @@ -684,16 +420,16 @@ int main() { float *fptr = sycl::malloc_shared(N, q); int *iptr = sycl::malloc_shared(N, q); - sycl::nd_range<1> ndr{N, WGSIZE}; + sycl::range<1> global{N}, local{WGSIZE}; // For a templated kernel, pass the address of a specific instantiation. - sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function>, 3.14f, fptr); - sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function>, 3, iptr); + sycl::khr::launch_grouped(q, global, local, sycl::khr::kernel_function>, 3.14f, fptr); + sycl::khr::launch_grouped(q, global, local, sycl::khr::kernel_function>, 3, iptr); // For an overloaded kernel, cast to the desired function type to select the // overload. - sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function<(void(*)(float*))ping>, fptr); - sycl::khr::nd_launch(q, ndr, sycl::khr::kernel_function<(void(*)(int*))ping>, iptr); + sycl::khr::launch_grouped(q, global, local, sycl::khr::kernel_function<(void(*)(float*))ping>, fptr); + sycl::khr::launch_grouped(q, global, local, sycl::khr::kernel_function<(void(*)(int*))ping>, iptr); q.wait(); From b460719721aea65a109043b80e0a8a827124ebce Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Tue, 7 Jul 2026 14:24:41 -0700 Subject: [PATCH 12/16] Simplify further, remove is_kernel_v trait, I considered it redundant --- .../sycl_khr_free_function_kernels.adoc | 46 +++---------------- 1 file changed, 7 insertions(+), 39 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 5140c5f25..2b943fd0e 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -3,8 +3,8 @@ This extension introduces _free function kernels_. A free function kernel is an ordinary C++ function, defined at namespace scope, -that is decorated with one of the [code]#SYCL_KHR_ND_KERNEL# or -[code]#SYCL_KHR_TASK_KERNEL# macros so that it becomes a device kernel entry +that is decorated with either the [code]#SYCL_KHR_ND_KERNEL# or +the [code]#SYCL_KHR_TASK_KERNEL# macro so that it becomes a device kernel entry point. Its kernel arguments are the function's parameters, rather than the captures of a lambda or the members of a function object, which gives them a defined order. @@ -15,13 +15,7 @@ a lambda. [[sec:khr-free-function-kernels-dependencies]] == Dependencies -This extension depends on the -<> extension. A free -function kernel obtains its work-item's position in the iteration space, when it -needs one, through the in-kernel queries provided by that extension (such as -[code]#sycl::khr::this_nd_item#). A kernel that does not need a position --- a -single-task free function kernel, or an ND-range free function kernel whose body -never queries one --- uses no such query. +This extension has no dependencies on other extensions. [[sec:khr-free-function-kernels-feature-test]] == Feature test macro @@ -43,8 +37,8 @@ below. == Defining a free function kernel A free function kernel is an ordinary C++ function whose declaration is decorated -with one of the [code]#SYCL_KHR_ND_KERNEL# or [code]#SYCL_KHR_TASK_KERNEL# -macros. +with either the [code]#SYCL_KHR_ND_KERNEL# or the [code]#SYCL_KHR_TASK_KERNEL# +macro. The kernel's _kind_ is determined by which macro is used: * [code]#SYCL_KHR_ND_KERNEL(N)# --- the function is launched over an @@ -93,9 +87,6 @@ A program that violates any of them is ill formed unless stated otherwise. an application can query this with the core [code]#sycl::is_device_copyable_v# type trait. -* No declaration of the function may specify a default argument for any - parameter. - * The decoration must appear on the first declaration of the function in the translation unit. A redeclaration of the function may also be decorated, provided it is decorated @@ -135,8 +126,8 @@ below are parameterized on it through a non-type template parameter == Traits for kernel functions This extension defines traits that report, at compile time, whether a given -[code]#Func# is a free function kernel (see -<>) and, if so, what kind it has. +[code]#Func# is decorated as a free function kernel of a particular kind (see +<>). In each trait, [code]#Func# is the address of a function, and the trait reports [code]#true# only when [code]#Func# is decorated as a free function kernel of the relevant kind. @@ -152,29 +143,6 @@ is unspecified.{endnote} ''' -.[apidef]#is_kernel# -[source,role=synopsis,id=api:khr-free-function-kernels-is_kernel] ----- -namespace sycl::khr { - -template -struct is_kernel; - -template -inline constexpr bool is_kernel_v = is_kernel::value; - -} // namespace sycl::khr ----- - -_Returns:_ [code]#is_kernel::value# is [code]#true# if [code]#Func# is the -address of a function that is decorated as a free function kernel (with any -kind), and [code]#false# otherwise. - -_Remarks:_ The expression [code]#is_kernel_v# is usable in a constant -expression. - -''' - .[apidef]#is_nd_kernel# [source,role=synopsis,id=api:khr-free-function-kernels-is_nd_kernel] ---- From 73f6e6f2818b32cc5c709218799b9dd9dac8a677 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Wed, 8 Jul 2026 07:25:38 -0700 Subject: [PATCH 13/16] Restrict free functions to not be sycl-external --- adoc/extensions/sycl_khr_free_function_kernels.adoc | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 2b943fd0e..11c3417b0 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -83,6 +83,11 @@ A program that violates any of them is ill formed unless stated otherwise. * The function must not accept a variadic argument list. +* The function must not be declared [code]#SYCL_EXTERNAL#. + A free function kernel is a kernel entry point, not a device function that is + called from device code, so the cross-translation-unit device linkage that + [code]#SYCL_EXTERNAL# provides does not apply to it. + * Each of the function's parameters must have a type that is <>; an application can query this with the core [code]#sycl::is_device_copyable_v# type trait. @@ -97,6 +102,7 @@ A program that violates any of them is ill formed unless stated otherwise. * If the function is decorated in one translation unit, every other translation unit that declares the same function must decorate it with the same macro (the same kind, and for [code]#SYCL_KHR_ND_KERNEL# the same [code]#N#). + A program that violates this rule is ill formed, no diagnostic is required. {note}Some SYCL types that are legal parameters for an ordinary SYCL kernel are not legal parameters for a free function kernel, because a free function kernel From d7c4e6d876d3575f6706cb7e6ab36df1c9da58c9 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 6 Aug 2026 10:41:09 -0700 Subject: [PATCH 14/16] Align presentation with content of PR --- .../sycl_khr_free_function_kernels.adoc | 163 ++++-------------- 1 file changed, 29 insertions(+), 134 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index 11c3417b0..d60d2c19f 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -3,13 +3,12 @@ This extension introduces _free function kernels_. A free function kernel is an ordinary C++ function, defined at namespace scope, -that is decorated with either the [code]#SYCL_KHR_ND_KERNEL# or -the [code]#SYCL_KHR_TASK_KERNEL# macro so that it becomes a device kernel entry -point. +that is decorated with the [code]#SYCL_KHR_KERNEL# macro so that it becomes +a device kernel entry point. Its kernel arguments are the function's parameters, rather than the captures of a lambda or the members of a function object, which gives them a defined order. Because it is a named function, the kernel is identified by the function itself, -so it can be launched and queried by reference to the function rather than only as +so it can be launched by reference to the function rather than only as a lambda. [[sec:khr-free-function-kernels-dependencies]] @@ -37,27 +36,14 @@ below. == Defining a free function kernel A free function kernel is an ordinary C++ function whose declaration is decorated -with either the [code]#SYCL_KHR_ND_KERNEL# or the [code]#SYCL_KHR_TASK_KERNEL# -macro. -The kernel's _kind_ is determined by which macro is used: +with the [code]#SYCL_KHR_KERNEL# macro. -* [code]#SYCL_KHR_ND_KERNEL(N)# --- the function is launched over an - [code]#nd_range# iteration space of [code]#N# dimensions; or - -* [code]#SYCL_KHR_TASK_KERNEL()# --- the function is launched as a single task, - without any iteration space. - -Neither macro takes any arguments beyond those shown above, though a future -extension may define more. +The macro takes no arguments, though a future extension may define more. [source,role=synopsis,id=api:khr-free-function-kernels-macro] ---- -// Decorate a function declaration to define an ND-range free function kernel -// over N dimensions. -#define SYCL_KHR_ND_KERNEL(N) /* see below */ - -// Decorate a function declaration to define a single-task free function kernel. -#define SYCL_KHR_TASK_KERNEL() /* see below */ +// Decorate a function declaration to define a free function kernel. +#define SYCL_KHR_KERNEL() ---- The first declaration of the function in the translation unit must be the @@ -66,7 +52,7 @@ For example: [source] ---- -SYCL_KHR_ND_KERNEL(1) +SYCL_KHR_KERNEL() void scale(float factor, float *data) { // ... } @@ -75,7 +61,7 @@ void scale(float factor, float *data) { A free function kernel must obey all of the following rules. A program that violates any of them is ill formed unless stated otherwise. -* The function must be decorated with exactly one of the two macros. +* The function must be decorated with the macro. * The function must be declared at namespace scope. @@ -93,15 +79,10 @@ A program that violates any of them is ill formed unless stated otherwise. [code]#sycl::is_device_copyable_v# type trait. * The decoration must appear on the first declaration of the function in the - translation unit. - A redeclaration of the function may also be decorated, provided it is decorated - with the same macro (the same kind, and for [code]#SYCL_KHR_ND_KERNEL# the same - [code]#N#). - The effect is the same whether or not a redeclaration is decorated. + translation unit. A redeclaration of the function must be decorated as well. * If the function is decorated in one translation unit, every other translation - unit that declares the same function must decorate it with the same macro (the - same kind, and for [code]#SYCL_KHR_ND_KERNEL# the same [code]#N#). + unit that declares the same function must decorate it with the same macro. A program that violates this rule is ill formed, no diagnostic is required. {note}Some SYCL types that are legal parameters for an ordinary SYCL kernel are @@ -112,88 +93,14 @@ For example, [code]#accessor#, [code]#local_accessor#, image accessors, <> but are not device-copyable and so cannot be used as free function kernel parameters.{endnote} -The way the function's body obtains its position in the iteration space must be -consistent with the kind implied by the macro used: a function decorated with -[code]#SYCL_KHR_ND_KERNEL(N)# obtains its position for an [code]#nd_range# of -[code]#N# dimensions, while a function decorated with -[code]#SYCL_KHR_TASK_KERNEL()# executes as a single work-item, with no iteration -space, and must not query an iteration position. -A program that violates this rule is ill formed, no diagnostic is required. - -A function decorated with one of these macros is a kernel entry point. -Calling it as an ordinary function, in either host or device code, results in +A function decorated with the [code]#SYCL_KHR_KERNEL# macro is a kernel entry point. +Calling the function as an ordinary function, in either host or device code, results in undefined behavior. -The function itself identifies the kernel: the launch and query APIs described +The function itself identifies the kernel: the launch APIs described below are parameterized on it through a non-type template parameter [code]#Func# (for example, [code]#sycl::khr::kernel_function#). -[[sec:khr-free-function-kernels-traits]] -== Traits for kernel functions - -This extension defines traits that report, at compile time, whether a given -[code]#Func# is decorated as a free function kernel of a particular kind (see -<>). -In each trait, [code]#Func# is the address of a function, and the trait reports -[code]#true# only when [code]#Func# is decorated as a free function kernel of the -relevant kind. -These traits are usable in constant expressions and in SFINAE contexts, so an -application can branch on, or constrain a template against, the kind of a free -function kernel without launching it. - -{note}The kernel whose address is [code]#Func# need not be defined in the same -translation unit as the use of the trait; the decorated declaration (with -[code]#SYCL_KHR_ND_KERNEL# or [code]#SYCL_KHR_TASK_KERNEL#) is sufficient. -How an implementation makes the decoration observable in a constant expression -is unspecified.{endnote} - -''' - -.[apidef]#is_nd_kernel# -[source,role=synopsis,id=api:khr-free-function-kernels-is_nd_kernel] ----- -namespace sycl::khr { - -template -struct is_nd_kernel; - -template -inline constexpr bool is_nd_kernel_v = is_nd_kernel::value; - -} // namespace sycl::khr ----- - -_Returns:_ [code]#is_nd_kernel::value# is [code]#true# if -[code]#Func# is the address of a function that is decorated as an ND-range free -function kernel over [code]#Dims# dimensions (with -[code]#SYCL_KHR_ND_KERNEL(Dims)#), and [code]#false# otherwise. - -_Remarks:_ The expression [code]#is_nd_kernel_v# is usable in a -constant expression. - -''' - -.[apidef]#is_task_kernel# -[source,role=synopsis,id=api:khr-free-function-kernels-is_task_kernel] ----- -namespace sycl::khr { - -template -struct is_task_kernel; - -template -inline constexpr bool is_task_kernel_v = is_task_kernel::value; - -} // namespace sycl::khr ----- - -_Returns:_ [code]#is_task_kernel::value# is [code]#true# if -[code]#Func# is the address of a function that is decorated as a single-task free -function kernel (with [code]#SYCL_KHR_TASK_KERNEL#), and [code]#false# otherwise. - -_Remarks:_ The expression [code]#is_task_kernel_v# is usable in a -constant expression. - [[sec:khr-free-function-kernels-launch]] == Launching a free function kernel @@ -246,14 +153,13 @@ void launch_task(handler &h, kernel_function_s k, Args&&... args); } // namespace sycl::khr ---- -_Constraints:_ Available only if [code]#is_task_kernel_v# is -[code]#true# and [code]#std::is_invocable_v# is -[code]#true#. +_Constraints:_ Available only if [code]#std::is_invocable_v# is +[code]#true# and [code]#Func# is decorated with the [code]#SYCL_KHR_KERNEL# macro. _Effects:_ Enqueues the free function kernel [code]#Func# to the queue [code]#q# (or to the command group of the handler [code]#h#) as a single task. Each value in the [code]#args# pack is passed to the corresponding parameter of -[code]#Func#, converting it to the parameter's type if necessary. +[code]#Func#. [[sec:khr-free-function-kernels-launch-launch_grouped]] === launch_grouped @@ -276,16 +182,15 @@ void launch_grouped(handler &h, range global, range local, } // namespace sycl::khr ---- -_Constraints:_ Available only if [code]#is_nd_kernel_v# is -[code]#true# and [code]#std::is_invocable_v# is -[code]#true#. +_Constraints:_ Available only if [code]#std::is_invocable_v# is +[code]#true# and [code]#Func# is decorated with the [code]#SYCL_KHR_KERNEL# macro. _Effects:_ Enqueues the free function kernel [code]#Func# to the queue [code]#q# (or to the command group of the handler [code]#h#) as an ND-range kernel, over the iteration space defined by the global range [code]#global# with work-group (local) size [code]#local#. Each value in the [code]#args# pack is passed to the corresponding parameter of -[code]#Func#, converting it to the parameter's type if necessary. +[code]#Func#. For example: @@ -294,25 +199,15 @@ For example: sycl::khr::launch_grouped(q, global, local, sycl::khr::kernel_function, factor, data); ---- -The [code]#Constraints# above check, at compile time, only the kernel's _kind_ -and _dimensionality_ and that the supplied arguments can be passed to -[code]#Func#. -They do not check the required work-group size of the kernel, if any, because -that value is compared against the local range, which is generally a run-time -value. - -A free function kernel is subject to the same launch-time requirements as any -other SYCL kernel (<>). -In particular, launching a kernel that requires a specific work-group size with a -local range that does not match that requirement throws a synchronous -[code]#exception# with the [code]#errc::nd_range# error code, exactly as for a -SYCL kernel decorated with [code]#reqd_work_group_size#. +The [code]#Constraints# above check, at compile time, that arguments can be passed to +[code]#Func#. They do not check the required work-group size of the kernel, +if any; that is checked at launch against the local range. [[sec:khr-free-function-kernels-example]] == Example The example below defines a free function kernel at namespace scope, decorates it -with [code]#SYCL_KHR_ND_KERNEL#, and launches it with +with [code]#SYCL_KHR_KERNEL#, and launches it with [code]#sycl::khr::launch_grouped#. The kernel obtains its work-item position with the <> query @@ -327,8 +222,8 @@ parameter. constexpr size_t N = 1024; constexpr size_t WGSIZE = 32; -// A free function kernel: an ordinary function, decorated with the kind. -SYCL_KHR_ND_KERNEL(1) +// A free function kernel: an ordinary function, decorated with the macro. +SYCL_KHR_KERNEL() void scale(float factor, float *data) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); data[i] *= factor; @@ -370,20 +265,20 @@ constexpr size_t WGSIZE = 32; // A templated free function kernel. template -SYCL_KHR_ND_KERNEL(1) +SYCL_KHR_KERNEL() void iota(T start, T *p) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); p[i] = start + static_cast(i); } // Two overloads of a free function kernel. -SYCL_KHR_ND_KERNEL(1) +SYCL_KHR_KERNEL() void ping(float *p) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); p[i] = 1.0f; } -SYCL_KHR_ND_KERNEL(1) +SYCL_KHR_KERNEL() void ping(int *p) { size_t i = sycl::khr::this_nd_item<1>().get_global_linear_id(); p[i] = 1; From 38047d842d9a58bdf135bb9b7adc006a7f3954fc Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 13 Aug 2026 09:00:16 -0700 Subject: [PATCH 15/16] [KHR][FFK] Require body position queries to match how the kernel is launched --- adoc/extensions/sycl_khr_free_function_kernels.adoc | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index d60d2c19f..c7fada7ad 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -97,6 +97,16 @@ A function decorated with the [code]#SYCL_KHR_KERNEL# macro is a kernel entry po Calling the function as an ordinary function, in either host or device code, results in undefined behavior. +The way a free function kernel's body obtains its work-item position must be +consistent with how the kernel is launched (see +<>). +A kernel launched as a single task with [code]#launch_task# executes as a single +work-item, with no iteration space, and must not query an iteration position. +A kernel launched over an [code]#nd_range# of [code]#Dims# dimensions with +[code]#launch_grouped# obtains its position for that iteration space, for example +through [code]#sycl::khr::this_nd_item#. +A program that violates this rule is ill formed, no diagnostic is required. + The function itself identifies the kernel: the launch APIs described below are parameterized on it through a non-type template parameter [code]#Func# (for example, [code]#sycl::khr::kernel_function#). From 74ca839af41cde99c380338e7d68da5e4720d7a1 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 13 Aug 2026 09:39:40 -0700 Subject: [PATCH 16/16] [KHR][FFK] Replace normative body-vs-launch rule with a non-normative note --- .../extensions/sycl_khr_free_function_kernels.adoc | 14 +++++--------- 1 file changed, 5 insertions(+), 9 deletions(-) diff --git a/adoc/extensions/sycl_khr_free_function_kernels.adoc b/adoc/extensions/sycl_khr_free_function_kernels.adoc index c7fada7ad..f695b6c3a 100644 --- a/adoc/extensions/sycl_khr_free_function_kernels.adoc +++ b/adoc/extensions/sycl_khr_free_function_kernels.adoc @@ -97,15 +97,11 @@ A function decorated with the [code]#SYCL_KHR_KERNEL# macro is a kernel entry po Calling the function as an ordinary function, in either host or device code, results in undefined behavior. -The way a free function kernel's body obtains its work-item position must be -consistent with how the kernel is launched (see -<>). -A kernel launched as a single task with [code]#launch_task# executes as a single -work-item, with no iteration space, and must not query an iteration position. -A kernel launched over an [code]#nd_range# of [code]#Dims# dimensions with -[code]#launch_grouped# obtains its position for that iteration space, for example -through [code]#sycl::khr::this_nd_item#. -A program that violates this rule is ill formed, no diagnostic is required. +{note}A free function kernel obtains its work-item position through the queries +provided by <>, such as +[code]#sycl::khr::this_nd_item#. The preconditions of those queries require the +query's dimensionality to match the dimensionality with which the kernel is +launched. See that extension for details.{endnote} The function itself identifies the kernel: the launch APIs described below are parameterized on it through a non-type template parameter