Skip to content

Commit fb626b6

Browse files
ProGTXProGTX
committed
CP024 Default placeholders
1. Deprecate `access::placeholder` 2. Extend placeholder functionality to ignore template parameter * Still applies only to global and constant buffers 3. Updated `handler::require` 4. New accessor constructors * Default constructor * From a buffer, 0-dim * From a buffer, ranged 5. New access member functions * `is_null` * `has_handler` 6. Vector addition example
1 parent 0a6be2a commit fb626b6

File tree

2 files changed

+270
-0
lines changed

2 files changed

+270
-0
lines changed

‎README.md

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -63,3 +63,4 @@ from this registry in the future.
6363
| CP020 | [Interop Task](interop_task/interop_task.md) | SYCL 1.2.1 | 16 January 2019 | 16 January 2019 | _Available since CE 1.0.5_ |
6464
| CP021 | [Default-Constructed Buffers](default-constructed-buffers/default-constructed-buffers.md) | SYCL 1.2.1 | 27 August 2019 | 5 September 2019 | _Draft_ |
6565
| CP022 | [Host Task with Interop capabilities](host_task/host_task.md) | SYCL 1.2.1 | 16 January 2019 | 20 January 2020 | _Final Draft_ |
66+
| CP024 | [Default placeholder](default-placeholders/index.md) | SYCL Next (after 1.2.1) | 9 March 2020 | 9 March 2020 | _Work in Progress_ |

‎default-placeholders/index.md

Lines changed: 269 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,269 @@
1+
# Default placeholders
2+
3+
| Proposal ID | CP024 |
4+
|-------------|--------|
5+
| Name | Default placeholders |
6+
| Date of Creation | 9 March 2019 |
7+
| Revision | 0.1 |
8+
| Latest Update | 9 March 2020 |
9+
| Target | SYCL Next (after 1.2.1) |
10+
| Current Status | _Work in Progress_ |
11+
| Reply-to | Peter Žužek <[email protected]> |
12+
| Original author | Peter Žužek <[email protected]> |
13+
| Contributors | Gordon Brown <[email protected]> |
14+
15+
## Overview
16+
17+
This proposal aims to deprecate `access::placeholder`
18+
and instead allow all accessors to global and constant memory to be placeholders.
19+
20+
## Revisions
21+
22+
### 0.1
23+
24+
* Initial proposal
25+
26+
## Motivation
27+
28+
SYCL 1.2.1 introduced the `access::placeholder` enumeration
29+
which is used as the 5th accessor template parameter
30+
to indicate whether the accessor can be used as a placeholder.
31+
Only the `global_buffer` and `constant_buffer` access targets
32+
support placeholder accessors.
33+
34+
The main reason for having placeholders is to store accessors into objects
35+
without having a queue at the point of object creation,
36+
and registering the access with a command group submission at a later time.
37+
One of Codeplay's proposals that didn't make it into SYCL 1.2.1
38+
was to allow placeholders to be default constructible as well,
39+
i.e. not having to even know the buffer at the point of construction.
40+
This extension is used in some SYCL ecosystem projects.
41+
42+
Accessors are pointer-like objects,
43+
and placeholders try to fill that gap
44+
that prevents accessors from being even more pointer-like.
45+
A default constructed placeholder accessor is analogous to a null pointer.
46+
A placeholder that's bound to a buffer
47+
but hasn't been registered with a command group
48+
is more like a fancy pointer:
49+
the user doesn't own the data
50+
until the accessor is registered and used in a kernel,
51+
where is becomes more similar to a regular pointer.
52+
53+
Having this type separation between full accessors and placeholders
54+
might be useful from a type safety perspective,
55+
but we believe it makes development more difficult.
56+
For example, [another one of our proposals](https://github.com/codeplaysoftware/standards-proposals/pull/100/files)
57+
introduces alias templates for different access targets
58+
and revises rules on how read-only accessors are handled,
59+
all in the name of reducing accessor verbosity.
60+
The placeholder template parameter makes that much more difficult,
61+
meaning we either need to introduce another parameter to the alias template,
62+
making it a lot less useful,
63+
or simply ignore the reduction in verbosity for placeholder accessors.
64+
65+
## Changes
66+
67+
### Deprecate `access::placeholder`
68+
69+
Mark the `access::placeholder` enum class as deprecated,
70+
but keep it for backwards compatibility
71+
until it's eventually removed from a subsequent standard.
72+
73+
```cpp
74+
namespace access {
75+
...
76+
77+
enum class placeholder // deprecated
78+
{...};
79+
} // namespace access
80+
81+
template <typename dataT,
82+
int dimensions,
83+
access::mode accessMode,
84+
access::target accessTarget = access::target::global_buffer,
85+
access::placeholder isPlaceholder = access::placeholder::false_t>
86+
class accessor;
87+
```
88+
89+
### All accessors with a global or constant target can be placeholders
90+
91+
SYCL 1.2.1 allows `access::placeholder::true_t`
92+
only when the access target is `global_buffer` or `constant_buffer`.
93+
We propose that the same placeholder semantics
94+
still apply to only these two targets,
95+
just that the semantics and API is available
96+
regardless of the `isPlaceholder` template parameter.
97+
98+
### Accept all accessors in `handler::require`
99+
100+
At the moment the member function `handler::require` only accepts
101+
placeholder accessors.
102+
This should stay the same,
103+
but since the enum class is deprecated,
104+
the function needs to be extended:
105+
106+
```cpp
107+
class handler {
108+
public:
109+
...
110+
111+
// Adds isPlaceholder to existing SYCL 1.2.1 function
112+
template <typename dataT,
113+
int dimensions,
114+
access::mode accessMode,
115+
access::target accessTarget,
116+
access::placeholder isPlaceholder>
117+
void require(accessor<dataT,
118+
dimensions,
119+
accessMode,
120+
accessTarget,
121+
isPlaceholder>
122+
acc);
123+
};
124+
```
125+
126+
### Deprecate `is_placeholder`
127+
128+
The function `accessor::is_placeholder` doesn't make sense anymore,
129+
we propose deprecating it.
130+
131+
### New constructors
132+
133+
We propose adding new constructors to the accessors class
134+
to allow placeholder construction.
135+
136+
1. Default constructor - not part of SYCL 1.2.1,
137+
but there is a [Codeplay proposal](https://github.com/codeplaysoftware/standards-proposals/pull/89)
138+
for making that happen.
139+
1. 0-dim accessor constructor from a buffer -
140+
normally an accessor constructor requires a buffer and a handler,
141+
we propose making the handler optional.
142+
This is the same constructor currently allowed for host buffers.
143+
1. Constructor from a buffer -
144+
normally an accessor constructor requires a buffer and a handler,
145+
we propose making the handler optional.
146+
This is the same constructor currently allowed for host buffers.
147+
148+
```cpp
149+
template <typename dataT,
150+
int dimensions,
151+
access::mode accessMode,
152+
access::target accessTarget = access::target::global_buffer,
153+
access::placeholder isPlaceholder = access::placeholder::false_t>
154+
class accessor {
155+
public:
156+
...
157+
158+
// 1
159+
// Only available when ((accessTarget == access::target::global_buffer) ||
160+
// (accessTarget == access::target::constant_buffer))
161+
accessor() noexcept;
162+
163+
// 2
164+
// Only available when: ((accessTarget == access::target::global_buffer) ||
165+
// (accessTarget == access::target::constant_buffer) ||
166+
// (accessTarget == access::target::host_buffer)) &&
167+
// (dimensions == 0)
168+
accessor(buffer<dataT, 1> &bufferRef);
169+
170+
// 3
171+
// Only available when: ((accessTarget == access::target::global_buffer) ||
172+
// (accessTarget == access::target::constant_buffer) ||
173+
// (accessTarget == access::target::host_buffer)) &&
174+
// (dimensions > 0)
175+
accessor(buffer<dataT, dimensions> &bufferRef,
176+
range<dimensions> accessRange,
177+
id<dimensions> accessOffset = {});
178+
};
179+
```
180+
181+
### New `accessor` member functions
182+
183+
In order to query the accessor for its status,
184+
we propose new member functions to the `accessor class`:
185+
186+
1. `is_null` - returns `true` if the accessor has been default constructed,
187+
which is only possible with placeholders.
188+
Not having an associated buffer is analogous to a null pointer.
189+
1. `has_handler` - returns `true` if the accessor is associated
190+
with a command group `handler`.
191+
Will only be `false` with host accessors and placeholder accessors.
192+
This replaces the `is_placeholder` member function.
193+
194+
```cpp
195+
template <typename dataT,
196+
int dimensions,
197+
access::mode accessMode,
198+
access::target accessTarget = access::target::global_buffer,
199+
access::placeholder isPlaceholder = access::placeholder::false_t>
200+
class accessor {
201+
public:
202+
...
203+
204+
// 1
205+
bool is_null() const noexcept;
206+
207+
// 2
208+
bool has_handler() const noexcept;
209+
};
210+
```
211+
212+
## Examples
213+
214+
Simple vector addition:
215+
216+
```cpp
217+
std::vector<int> a{1, 2, 3, 4, 5};
218+
std::vector<int> b{6, 7, 8, 9, 10};
219+
220+
using read_acc =
221+
accessor<int, 1, access::mode::read, access::target::global_buffer>;
222+
using write_acc =
223+
accessor<int, 1, access::mode::discard_write, access::target::global_buffer>;
224+
225+
read_acc accA;
226+
read_acc accB;
227+
write_acc accC;
228+
229+
// Sanity checks
230+
assert(accA.is_null());
231+
assert(!accA.has_handler());
232+
233+
const auto N = a.size();
234+
const auto bufRange = range<1>(N);
235+
236+
queue myQueue;
237+
238+
// Create a buffer and copy `a` into it
239+
buffer<int> bufA{bufRange};
240+
241+
accA = read_acc{bufA};
242+
assert(!accA.is_null());
243+
assert(!accA.has_handler());
244+
245+
myQueue.submit([&](handler &cgh) {
246+
cgh.require(accA);
247+
cgh.copy(a.data(), accA);
248+
});
249+
250+
// Create a buffer and copy `b` into it
251+
buffer<int> bufB{bufRange};
252+
accB = read_acc{bufB};
253+
myQueue.submit([&](handler &cgh) {
254+
cgh.require(accB);
255+
cgh.copy(b.data(), accB);
256+
});
257+
258+
// Submit kernel that writes to output buffer
259+
// Use constant buffer accessors
260+
buffer<int> bufC{bufRange};
261+
accC = read_acc{bufC};
262+
myQueue.submit([&](handler &cgh) {
263+
cgh.require(accA);
264+
cgh.require(accB);
265+
cgh.require(accC);
266+
cgh.parallel_for<class vec_add>(bufRange,
267+
[=](id<1> i) { accC[i] = accA[i] + accB[i]; });
268+
});
269+
```

0 commit comments

Comments
 (0)