@@ -43,8 +43,7 @@ SYCL specification refer to that revision.
4343The following extensions are required:
4444
4545- link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties]
46-
47- - link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[sycl_ext_oneapi_kernel_properties]
46+ - link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions]
4847
4948
5049== Status
@@ -112,9 +111,9 @@ in device local memory.
112111
113112=== Kernel properties
114113
115- The `work_group_scratch_size` property must be passed to a kernel to determine
116- the run-time size of the device local memory allocation associated with
117- all `get_work_group_scratch_memory` calls.
114+ The `work_group_scratch_size` property must be passed as a kernel launch
115+ property to determine the run-time size of the device local memory allocation
116+ associated with all `get_work_group_scratch_memory` calls.
118117
119118[source,c++]
120119----
@@ -151,38 +150,33 @@ then the implementation must throw a synchronous exception with the
151150local memory into account, whether via the `work_group_scratch_size` property
152151or other APIs such as `local_accessor`.
153152
153+
154154== Example
155155
156156[source,c++]
157157----
158+ #include <sycl/sycl.hpp>
158159namespace syclex = sycl::ext::oneapi::experimental;
159160
160- ...
161-
162- q.parallel_for(sycl::nd_range<1>{N, M},
163- syclex::properties{syclex::work_group_scratch_size(M * sizeof(int))},
164- [=](sycl::nd_item<1> it) {
165- auto ptr = syclex::get_work_group_scratch_memory();
166- auto ptr2 = syclex::get_work_group_scratch_memory();
167- });
168- ----
169-
161+ constexpr size_t GLOBAL = 1024;
162+ constexpr size_t N = 256;
170163
171- == Implementation notes
164+ int main() {
165+ sycl::queue q;
172166
173- This non-normative section provides information about one possible
174- implementation of this extension. It is not part of the specification of the
175- extension's API.
167+ syclex::launch_config cfg{
168+ sycl::nd_range{{GLOBAL}, {N}},
169+ syclex::properties{
170+ syclex::work_group_scratch_size{N * sizeof(int)}
171+ }
172+ };
176173
177- For `get_work_group_scratch_memory`,
178- the implementation may need to generate some additional code to
179- appropriately initialize the pointer(s) returned by the call.
180- Alternatively, it may be possible to initialize the pointer to the beginning
181- of the device's local memory region (if that value is known). Either way, the
182- implementation must account for the existence of one or more `local_accessor`
183- objects (which themselves may allocate a dynamic amount of device local
184- memory).
185-
186-
187- == Issues
174+ syclex::nd_launch(q, cfg, [=](sycl::nd_item<> it) {
175+ // Gets a pointer to N int's
176+ int *ptr = static_cast<int *>(syclex::get_work_group_scratch_memory());
188177
178+ // Gets the same pointer to N int's
179+ int *ptr2 = static_cast<int *>(syclex::get_work_group_scratch_memory());
180+ });
181+ }
182+ ----
0 commit comments