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..b05997e15a6cf 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -287,16 +287,17 @@ 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 -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 { @@ -363,6 +364,87 @@ 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 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 { + +template <typename Function, typename Properties = empty_properties_t> +struct kernel_function { + + kernel_function(Function &&f, Properties p = properties{}); + + // Available only if Function is invocable with Args + template <typename... Args> + void operator()(Args&&... args) const; + + // Available only if Properties contains no run-time properties + static constexpr auto get(properties_tag) const; + + // Available only if Properties contains at least one run-time property + auto get(properties_tag) const; + +} // namespace sycl::ext::oneapi::experimental +``` + +--- + +```c++ +template <typename... Args> +void operator()(Args&&... args) const; +``` + +_Constraints_: `Function` is invocable with `Args`. + +_Effects_: Invokes `Function` with `Args`. + +--- + +```c++ +static constexpr auto get(properties_tag) const; (1) + +auto get(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: + +```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