From f6581d6612788e21d9354afc3b14ec7eab25f9db Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 25 Mar 2025 11:20:58 +0000 Subject: [PATCH 1/6] [SYCL][Doc] Add kernel_function lambda wrapper We previously removed the ability to pass properties directly to functions like single_task, parallel_for, etc, in favor of requiring them to be attached to the kernel. By introducing a new kernel_function wrapper that attaches properties to a kernel function, we remove the need for (most) developers to write such wrappers themselves. Signed-off-by: John Pennycook --- ...sycl_ext_oneapi_kernel_properties.asciidoc | 58 ++++++++++++++++++- 1 file changed, 57 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc index 2ca9ac1b55d2a..42685228c1eaf 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -287,7 +287,7 @@ by the property, the implementation must throw a synchronous exception with the |=== -== Embedding Properties into a Kernel +=== Embedding Properties into a Kernel In other situations it may be useful to embed a kernel's properties directly into its type, to ensure that a kernel cannot be launched without a property @@ -363,6 +363,62 @@ diagnostic; invalid combinations that can only be detected at run-time should result in an implementation throwing an `exception` with the `errc::invalid` error code. +=== Using Properties with Lambda Expressions + +When a SYCL kernel is defined via a lambda expression, there is no way to +define a `get` member function and subsequently no way to embed kernel +properties. Instead, developers must wrap the lambda expression in an object. + +To simplify this usage pattern, this extension defines a `kernel_function` +that encapsulates a kernel function (which may be a lambda expression) and a +property list. + +NOTE: Developers are free to extend `kernel_function` or define their own +wrapper classes (e.g., to attach commonly used property lists). + +```c++ +namespace sycl::ext::oneapi::experimental { + +template +struct kernel_function { + + kernel_function(Function f, Properties p = syclx::properties{}); + + // Available only if Function is invocable with Args + template + void operator()(Args... args) const { + f(std::forward(args)...); + } + + auto get(syclx::properties_tag) { + return Properties{}; + } + + private: + const Function f; // exposition only + +} // namespace sycl::ext::oneapi::experimental +``` + +The example below shows how the `KernelFunctor` example from the previous +section can be written using this wrapper: + +```c++ +namespace syclx = sycl::ext::oneapi::experimental; + +... + +auto lambda = [=](id<1> i) const { + a[i] = b[i] + c[i]; +} +auto props = syclx::properties{syclx::work_group_size<8, 8>, syclx::sub_group_size<8>}; +auto kernel = syclx::kernel_function(lambda, props); + +... + +q.parallel_for(range<2>{16, 16}, kernel).wait(); +``` + === Querying Properties in a Compiled Kernel Any properties embedded into a kernel type via a property list are reflected From 82de663d31337eed1f70e66819b2ffe229d9de0f Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 25 Mar 2025 11:26:21 +0000 Subject: [PATCH 2/6] Fix a few references to old behavior While working on the kernel_function definition, I noticed that some text referred to "previous sections" that no longer exist. I tweaked the wording here to reflect that properties must always be attached to the kernel type. Signed-off-by: John Pennycook --- ...sycl_ext_oneapi_kernel_properties.asciidoc | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc index 42685228c1eaf..eb934833261be 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -289,14 +289,15 @@ by the property, the implementation must throw a synchronous exception with the === Embedding Properties into a Kernel -In other situations it may be useful to embed a kernel's properties directly -into its type, to ensure that a kernel cannot be launched without a property -that it depends upon for correctness. +A kernel's properties are embedded directly into its type, to ensure that a +kernel cannot be launched without a property that it depends upon for +correctness. -To enable this use-case, this extension adds a mechanism for implementations to -extract a property list from a kernel functor, if a kernel functor declares -a member function named `get` accepting a `sycl::ext::oneapi::experimental::properties_tag` -tag type and returning an instance of `sycl::ext::oneapi::experimental::properties`. +To enable this, this extension adds a mechanism for implementations to extract +a property list from a kernel functor, if a kernel functor declares a member +function named `get` accepting a +`sycl::ext::oneapi::experimental::properties_tag` tag type and returning an +instance of `sycl::ext::oneapi::experimental::properties`. ```c++ namespace sycl { @@ -323,8 +324,8 @@ attributes to be applied to different call operators within the same functor. An embedded property list applies to all call operators in the functor. -The example below shows how the kernel from the previous section could be -rewritten to leverage an embedded property list: +The example below shows how a simple vector addition kernel could be +written to leverage an embedded property list: ```c++ struct KernelFunctor { From be618afbf08b82b876c2e4768a5261141091e31f Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 25 Mar 2025 15:03:09 +0000 Subject: [PATCH 3/6] Remove implementation details from synopsis Replace this with a detailed description of what the functions do. This is intended to permit implementations to specialize when properties is empty, to avoid transferring properties objects unnecessarily. Signed-off-by: John Pennycook --- ...sycl_ext_oneapi_kernel_properties.asciidoc | 40 +++++++++++++++---- 1 file changed, 32 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc index eb934833261be..639d2ec72bf4c 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -387,20 +387,44 @@ struct kernel_function { // Available only if Function is invocable with Args template - void operator()(Args... args) const { - f(std::forward(args)...); - } + void operator()(Args... args) const; - auto get(syclx::properties_tag) { - return Properties{}; - } + // Available only if Properties contains no run-time properties + static constexpr auto get(syclx::properties_tag) const; - private: - const Function f; // exposition only + // Available only if Properties contains at least one run-time property + auto get(syclx::properties_tag) const; } // namespace sycl::ext::oneapi::experimental ``` +--- + +```c++ +template +void operator()(Args... args) const; +``` + +_Constraints_: `Function` is invocable with `Args`. + +_Effects_: Invokes `Function` with `Args`. + +--- + +```c++ +static constexpr auto get(syclx::properties_tag) const; (1) + +auto get(syclx::properties_tag) const; (2) +``` + +_Constraints_ (1): `Properties` contains no run-time properties. + +_Constraints_ (2): `Properties` contains at least one run-time property. + +_Returns_: The property list associated with this kernel function. + +--- + The example below shows how the `KernelFunctor` example from the previous section can be written using this wrapper: From 3d580ff5d62745c227c3c323ca4dc04483961bad Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 25 Mar 2025 15:11:42 +0000 Subject: [PATCH 4/6] Add && to Function and Args Signed-off-by: John Pennycook --- .../experimental/sycl_ext_oneapi_kernel_properties.asciidoc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc index 639d2ec72bf4c..d0af766d7aa47 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -383,11 +383,11 @@ namespace sycl::ext::oneapi::experimental { template struct kernel_function { - kernel_function(Function f, Properties p = syclx::properties{}); + kernel_function(Function &&f, Properties p = syclx::properties{}); // Available only if Function is invocable with Args template - void operator()(Args... args) const; + void operator()(Args&&... args) const; // Available only if Properties contains no run-time properties static constexpr auto get(syclx::properties_tag) const; @@ -402,7 +402,7 @@ struct kernel_function { ```c++ template -void operator()(Args... args) const; +void operator()(Args&&... args) const; ``` _Constraints_: `Function` is invocable with `Args`. From 160fff74626f502bf3feca86bdc556d519327d42 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 25 Mar 2025 15:19:11 +0000 Subject: [PATCH 5/6] Remove syclx:: from synopsis and definitions syclx:: convention is only for examples. Signed-off-by: John Pennycook --- .../sycl_ext_oneapi_kernel_properties.asciidoc | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc index d0af766d7aa47..fd4e80870721e 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -383,17 +383,17 @@ namespace sycl::ext::oneapi::experimental { template struct kernel_function { - kernel_function(Function &&f, Properties p = syclx::properties{}); + kernel_function(Function &&f, Properties p = properties{}); // Available only if Function is invocable with Args template void operator()(Args&&... args) const; // Available only if Properties contains no run-time properties - static constexpr auto get(syclx::properties_tag) const; + static constexpr auto get(properties_tag) const; // Available only if Properties contains at least one run-time property - auto get(syclx::properties_tag) const; + auto get(properties_tag) const; } // namespace sycl::ext::oneapi::experimental ``` @@ -412,9 +412,9 @@ _Effects_: Invokes `Function` with `Args`. --- ```c++ -static constexpr auto get(syclx::properties_tag) const; (1) +static constexpr auto get(properties_tag) const; (1) -auto get(syclx::properties_tag) const; (2) +auto get(properties_tag) const; (2) ``` _Constraints_ (1): `Properties` contains no run-time properties. From 174354294c3fda6aee013c5bddc3da7925473096 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 4 Apr 2025 07:48:07 +0100 Subject: [PATCH 6/6] Clarify meaning of "extend" kernel_function Signed-off-by: John Pennycook --- .../experimental/sycl_ext_oneapi_kernel_properties.asciidoc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc index fd4e80870721e..b05997e15a6cf 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -374,8 +374,9 @@ To simplify this usage pattern, this extension defines a `kernel_function` that encapsulates a kernel function (which may be a lambda expression) and a property list. -NOTE: Developers are free to extend `kernel_function` or define their own -wrapper classes (e.g., to attach commonly used property lists). +NOTE: Developers are free to create classes that derive from `kernel_function` +or define their own wrapper classes (e.g., to attach commonly used property +lists). ```c++ namespace sycl::ext::oneapi::experimental {