@@ -120,29 +120,29 @@ using the `private_alloca` API defined in the following sections.
120
120
[source,c++]
121
121
----
122
122
namespace sycl::ext::oneapi::experimental {
123
- template <typename ElementType, auto &SpecName ,
123
+ template <typename ElementType, auto &SizeSpecName ,
124
124
access::decorated DecorateAddress>
125
125
private_ptr<ElementType, DecorateAddress>
126
126
private_alloca(kernel_handler &kh);
127
127
128
- template <typename ElementType, std::size_t Alignment, auto &SpecName ,
128
+ template <typename ElementType, std::size_t Alignment, auto &SizeSpecName ,
129
129
access::decorated DecorateAddress>
130
130
private_ptr<ElementType, DecorateAddress>
131
131
aligned_private_alloca(kernel_handler &kh);
132
132
} // namespace sycl::ext::oneapi::experimental
133
133
----
134
134
135
- _Mandates_: `ElementType` must be a cv-unqualified trivial type and `SpecName`
136
- must be a reference to a specialization constant of integral `value_type`. In
137
- the case of `aligned_private_alloca`, `Alignment` must be an alignment value and
138
- must be a positive multiple of `alignof(ElementType)`. If `Alignment` is an
139
- extended alignment, it must be supported by the implementation.
135
+ _Mandates_: `ElementType` must be a cv-unqualified trivial type and
136
+ `SizeSpecName` must be a reference to a specialization constant of integral
137
+ type. In the case of `aligned_private_alloca`, `Alignment` must be an alignment
138
+ value and must be a positive multiple of `alignof(ElementType)`. If `Alignment`
139
+ is an extended alignment, it must be supported by the implementation.
140
140
141
- _Effects_: `h .get_specialization_constant<size >()` elements of type
141
+ _Effects_: `kh .get_specialization_constant<SizeSpecName >()` elements of type
142
142
`ElementType` are allocated and default initialized in private memory.
143
143
144
144
_Returns_: A pointer to a default initialized region of private memory of
145
- `h .get_specialization_constant<size >()` elements of type
145
+ `kh .get_specialization_constant<SizeSpecName >()` elements of type
146
146
`ElementType`. `DecorateAddress` defines whether the returned `multi_ptr` is
147
147
decorated. In the case of `private_alloca`, the pointer is suitably aligned for
148
148
an object of type `ElementType`. In the case of `aligned_private_alloca`, the
@@ -151,7 +151,7 @@ pointer is aligned to the specified `Alignment`.
151
151
_Remarks_: In case of private memory exhaustion, the implementation must report
152
152
an error in the same fashion as if the allocation size were static. In case of a
153
153
successful call, allocated memory has automatic storage duration. Additionally,
154
- `SpecName ` must have a default value of at least 1 and not be set to a value
154
+ `SizeSpecName ` must have a default value of at least 1 and not be set to a value
155
155
less than 1 during program execution. Violation of these conditions results in
156
156
undefined behaviour.
157
157
@@ -213,6 +213,35 @@ void run(queue q, const float *in, float *out, size_t n) {
213
213
});
214
214
----
215
215
216
+ === Usage with `sycl::span`
217
+
218
+ In this section, we show an example of how users could use this extension with
219
+ `sycl::span` as a `std::array` replacement:
220
+
221
+ [source,c++]
222
+ ----
223
+ constexpr specialization_id<std::size_t> size(1);
224
+
225
+ class Kernel;
226
+
227
+ // Counterpart to 'impl' in the first example using 'sycl::span'
228
+ SYCL_EXTERNAL void impl(const float *in, float *out,
229
+ sycl::span<float> ptr);
230
+
231
+ void run(queue q, const float *in, float *out, size_t n) {
232
+ q.submit([&](handler &h) {
233
+ h.set_specialization_constant<size>(n);
234
+ h.parallel_for<Kernel>(n, [=](id<1> i, kernel_handler kh) {
235
+ // Create sycl::span with the returned pointer and the specialization
236
+ // constant used as size.
237
+ sycl::span<float> tmp{
238
+ private_alloca<float, size, access::decorated::no>(kh).get_raw(),
239
+ kh.get_specialization_constant<size>()};
240
+ impl(in, out, tmp);
241
+ });
242
+ });
243
+ ----
244
+
216
245
== Design constraints
217
246
218
247
The big design constraint stems from the unknown allocation size at compile
0 commit comments