Skip to content

Commit f357e69

Browse files
authored
[SYCL][Doc] Add free function kernel spec example (#20507)
Add an example showing how to use sycl_ext_oneapi_work_group_scratch_memory from a free function kernel.
1 parent 965e4fb commit f357e69

File tree

1 file changed

+39
-0
lines changed

1 file changed

+39
-0
lines changed

sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -994,6 +994,45 @@ int main() {
994994
}
995995
----
996996

997+
=== Using "scratch" work-group local memory
998+
999+
Free function kernels can use work-group local memory via `local_accessor` or
1000+
via other extensions.
1001+
This example demonstrates how to use work-group local memory via the
1002+
link:../experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc[
1003+
sycl_ext_oneapi_work_group_scratch_memory] extension because it also illustrates
1004+
how to pass a kernel launch parameter.
1005+
1006+
[source,c++]
1007+
----
1008+
#include <sycl/sycl.hpp>
1009+
namespace syclexp = sycl::ext::oneapi::experimental;
1010+
1011+
static constexpr size_t NUM = 1024;
1012+
static constexpr size_t WGSIZE = 16;
1013+
1014+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
1015+
void mykernel() {
1016+
// Gets a pointer to WGSIZE int's
1017+
int *ptr = static_cast<int *>(syclexp::get_work_group_scratch_memory());
1018+
}
1019+
1020+
int main() {
1021+
sycl::queue q;
1022+
1023+
syclexp::launch_config cfg{
1024+
sycl::nd_range{{NUM}, {WGSIZE}},
1025+
syclexp::properties{
1026+
syclexp::work_group_scratch_size{WGSIZE * sizeof(int)}
1027+
}
1028+
};
1029+
syclexp::nd_launch(q, cfg, syclexp::kernel_function<mykernel>);
1030+
1031+
q.wait();
1032+
}
1033+
----
1034+
1035+
9971036

9981037
[[level-zero-and-opencl-compatibility]]
9991038
== {dpcpp} guaranteed compatibility with Level Zero and OpenCL backends

0 commit comments

Comments
 (0)