Skip to content

Commit acf2ca8

Browse files
authored
[SYCL][Doc] Add indeterminate to work_group_memory (#15933)
We decided that it was too easy to mistakenly use the default constructor like this, with the expectation that it statically allocates work-group local memory: ``` void device_code() { syclex::work_group_memory<int> mem; } ``` To make this error less likely, we add a parameter to the constructor named `indeterminate`: ``` void device_code() { syclex::work_group_memory<int> mem{syclex::indeterminate}; } ``` We hope this will make it more apparent that `mem` is just a dummy object, and it needs to be assigned to some other `work_group_memory` object before it can be used.
1 parent 4444135 commit acf2ca8

File tree

1 file changed

+11
-3
lines changed

1 file changed

+11
-3
lines changed

sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_memory.asciidoc

Lines changed: 11 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -106,12 +106,15 @@ This extension adds the following new class:
106106
----
107107
namespace sycl::ext::oneapi::experimental {
108108
109+
struct indeterminate_t {};
110+
inline constexpr indeterminate_t indeterminate;
111+
109112
template<typename DataT, typename PropertyListT = empty_properties_t>
110113
class work_group_memory {
111114
public:
112115
using value_type = std::remove_all_extents_t<DataT>;
113116
114-
work_group_memory();
117+
work_group_memory(const indeterminate_t&);
115118
work_group_memory(const work_group_memory& rhs);
116119
work_group_memory(handler& cgh);
117120
work_group_memory(size_t num, handler& cgh);
@@ -132,7 +135,7 @@ The `work_group_memory` class allocates device local memory and provides access
132135
to this memory from within a SYCL kernel function.
133136
The local memory that is allocated is shared between all work-items of a
134137
work-group.
135-
If multiple work-groups execute simultaneously, each of those work-group
138+
If multiple work-groups execute simultaneously, each of those work-groups
136139
receives its own independent copy of the allocated local memory.
137140

138141
The `work_group_memory` type is a legal kernel parameter type as defined in
@@ -193,7 +196,7 @@ array extents removed.
193196
a@
194197
[source,c++]
195198
----
196-
work_group_memory();
199+
work_group_memory(const indeterminate_t&);
197200
----
198201
!====
199202

@@ -207,6 +210,11 @@ member functions or operators produces undefined behavior.
207210
[_Note:_ This constructor may be called in either host code or device code.
208211
_{endnote}_]
209212

213+
[_Note:_ The parameter of type `indeterminate_t` is present only to help make it
214+
apparent that this constructor creates a "dummy" `work_group_memory` object.
215+
Applications typically pass the value `indeterminate`.
216+
_{endnote}_]
217+
210218
'''
211219

212220
[frame=all,grid=none,separator="@"]

0 commit comments

Comments
 (0)