From 01f298a90da3ac352f3f8b91a9eb551544cd682e Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Wed, 29 Oct 2025 15:58:14 -0400 Subject: [PATCH 1/2] [SYCL][Doc] Fix work_group_scratch_memory example The example in the specification was using an old API that we removed in #14785. Update the example and the spec wording to use `launch_config` from sycl_ext_oneapi_enqueue_functions to pass the property with the size of the scratch memory. --- ..._oneapi_work_group_scratch_memory.asciidoc | 56 +++++++++---------- 1 file changed, 25 insertions(+), 31 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc index 757406b068fa4..2ec36770c853d 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc @@ -36,15 +36,14 @@ https://github.com/intel/llvm/issues == Dependencies -This extension is written against the SYCL 2020 revision 9 specification. All +This extension is written against the SYCL 2020 revision 10 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision. The following extensions are required: - link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] - -- link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[sycl_ext_oneapi_kernel_properties] +- link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions] == Status @@ -112,9 +111,9 @@ in device local memory. === Kernel properties -The `work_group_scratch_size` property must be passed to a kernel to determine -the run-time size of the device local memory allocation associated with -all `get_work_group_scratch_memory` calls. +The `work_group_scratch_size` property must be passed as a kernel launch +property to determine the run-time size of the device local memory allocation +associated with all `get_work_group_scratch_memory` calls. [source,c++] ---- @@ -151,38 +150,33 @@ then the implementation must throw a synchronous exception with the local memory into account, whether via the `work_group_scratch_size` property or other APIs such as `local_accessor`. + == Example [source,c++] ---- +#include namespace syclex = sycl::ext::oneapi::experimental; -... - -q.parallel_for(sycl::nd_range<1>{N, M}, - syclex::properties{syclex::work_group_scratch_size(M * sizeof(int))}, - [=](sycl::nd_item<1> it) { - auto ptr = syclex::get_work_group_scratch_memory(); - auto ptr2 = syclex::get_work_group_scratch_memory(); -}); ----- - +constexpr size_t GLOBAL = 1024; +constexpr size_t N = 256; -== Implementation notes +int main() { + sycl::queue q; -This non-normative section provides information about one possible -implementation of this extension. It is not part of the specification of the -extension's API. + syclex::launch_config cfg{ + sycl::nd_range{{GLOBAL}, {N}}, + syclex::properties{ + syclex::work_group_scratch_size{N * sizeof(int)} + } + }; -For `get_work_group_scratch_memory`, -the implementation may need to generate some additional code to -appropriately initialize the pointer(s) returned by the call. -Alternatively, it may be possible to initialize the pointer to the beginning -of the device's local memory region (if that value is known). Either way, the -implementation must account for the existence of one or more `local_accessor` -objects (which themselves may allocate a dynamic amount of device local -memory). - - -== Issues + syclex::nd_launch(q, cfg, [=](sycl::nd_item<> it) { + // Gets a pointer to N int's + int *ptr = static_cast(syclex::get_work_group_scratch_memory()); + // Gets the same pointer to N int's + int *ptr2 = static_cast(syclex::get_work_group_scratch_memory()); + }); +} +---- From 1cf4ff9af73728af8c4e96b019f2c9639de9a1fe Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 14 Nov 2025 07:49:25 -0500 Subject: [PATCH 2/2] Remove "constexpr" from example These `constexpr` declarations aren't necessary, and they make the example more confusing. The purpose of this extension is to allocate a runtime-defined amount of local memory, so the memory size shouldn't be declared `constexpr`. --- .../sycl_ext_oneapi_work_group_scratch_memory.asciidoc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc index 2ec36770c853d..fcc7247b9289a 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc @@ -36,7 +36,7 @@ https://github.com/intel/llvm/issues == Dependencies -This extension is written against the SYCL 2020 revision 10 specification. All +This extension is written against the SYCL 2020 revision 11 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision. @@ -158,8 +158,8 @@ or other APIs such as `local_accessor`. #include namespace syclex = sycl::ext::oneapi::experimental; -constexpr size_t GLOBAL = 1024; -constexpr size_t N = 256; +size_t GLOBAL = 1024; +size_t N = 256; int main() { sycl::queue q;