Skip to content

[SYCL][Doc] Fix work_group_scratch_memory example #20506

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 2 commits into
base: sycl
Choose a base branch
from
Open

[SYCL][Doc] Fix work_group_scratch_memory example #20506

Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
01f298a
[SYCL][Doc] Fix work_group_scratch_memory example
gmlueck Oct 29, 2025
1cf4ff9
Remove "constexpr" from example
gmlueck Nov 14, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
56 changes: 25 additions & 31 deletions .../doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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++]
----
Expand Down Expand Up @@ -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 <sycl/sycl.hpp>
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<int *>(syclex::get_work_group_scratch_memory());

// Gets the same pointer to N int's
int *ptr2 = static_cast<int *>(syclex::get_work_group_scratch_memory());
});
}
----