diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_memory.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_memory.asciidoc index c142e66ec4c77..296b77acf82fb 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_memory.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_memory.asciidoc @@ -106,12 +106,15 @@ This extension adds the following new class: ---- namespace sycl::ext::oneapi::experimental { +struct indeterminate_t {}; +inline constexpr indeterminate_t indeterminate; + template class work_group_memory { public: using value_type = std::remove_all_extents_t; - work_group_memory(); + work_group_memory(const indeterminate_t&); work_group_memory(const work_group_memory& rhs); work_group_memory(handler& cgh); work_group_memory(size_t num, handler& cgh); @@ -132,7 +135,7 @@ The `work_group_memory` class allocates device local memory and provides access to this memory from within a SYCL kernel function. The local memory that is allocated is shared between all work-items of a work-group. -If multiple work-groups execute simultaneously, each of those work-group +If multiple work-groups execute simultaneously, each of those work-groups receives its own independent copy of the allocated local memory. The `work_group_memory` type is a legal kernel parameter type as defined in @@ -193,7 +196,7 @@ array extents removed. a@ [source,c++] ---- -work_group_memory(); +work_group_memory(const indeterminate_t&); ---- !==== @@ -207,6 +210,11 @@ member functions or operators produces undefined behavior. [_Note:_ This constructor may be called in either host code or device code. _{endnote}_] +[_Note:_ The parameter of type `indeterminate_t` is present only to help make it +apparent that this constructor creates a "dummy" `work_group_memory` object. +Applications typically pass the value `indeterminate`. +_{endnote}_] + ''' [frame=all,grid=none,separator="@"]