From 965538284927566d30928ccfc89d73d1aac529cb Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Wed, 29 Oct 2025 17:07:49 -0400 Subject: [PATCH] [SYCL][Doc] Add free function kernel spec example Add an example showing how to use sycl_ext_oneapi_work_group_scratch_memory from a free function kernel. --- ..._ext_oneapi_free_function_kernels.asciidoc | 40 +++++++++++++++++++ 1 file changed, 40 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc index 8827b4ca7e5f4..f949a2fa0142c 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc @@ -994,6 +994,46 @@ int main() { } ---- +=== Using "scratch" work-group local memory + +Free function kernels can use work-group local memory via `local_accessor` or +via other extensions. +This example demonstrates how to use work-group local memory via the +link:../experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc[ +sycl_ext_oneapi_work_group_scratch_memory] extension because it also illustrates +how to pass a kernel launch parameter. + +[source,c++] +---- +#include +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void mykernel() { + // Gets a pointer to WGSIZE int's + int *ptr = static_cast(syclexp::get_work_group_scratch_memory()); +} + +int main() { + sycl::queue q; + + syclexp::launch_config cfg{ + sycl::nd_range{{NUM}, {WGSIZE}}, + syclexp::properties{ + syclexp::work_group_scratch_size{WGSIZE * sizeof(int)} + } + }; + syclexp::nd_launch(q, cfg, syclexp::kernel_function); + + q.wait(); +} +---- + + [[level-zero-and-opencl-compatibility]] == {dpcpp} guaranteed compatibility with Level Zero and OpenCL backends