@@ -10,18 +10,16 @@ optional kernel features.
10
10
11
11
The requirements for this design come mostly from the SYCL 2020 specification
12
12
[ section 5.7 "Optional kernel features"] [ 1 ] but they also encompass the C++
13
- attribute ` [[sycl::requires ()]] ` that is described in [ section 5.8.1 "Kernel
14
- attributes"] [ 2 ] and [ section 5.8.2 "Device function attributes"] [ 3 ] .
13
+ attribute ` [[sycl::device_has ()]] ` that is described in section 5.8.1 "Kernel
14
+ attributes" and section 5.8.2 "Device function attributes".
15
15
16
16
[ 1 ] : < https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:optional-kernel-features >
17
- [ 2 ] : < https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.attributes >
18
- [ 3 ] : < https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_device_function_attributes >
19
17
20
- ** NOTE ** : At the time this document was written, there is a
21
- [ proposed change ] [ 4 ] to the SYCL 2020 specification that will rename
22
- ` [[sycl::requires()]] ` to ` [[sycl::device_has ()]] ` . Since that proposal has
23
- not yet been adopted, this design document continues to use the
24
- ` [[sycl::requires()]] ` name .
18
+ At the time this document was written, the published version of the SYCL 2020
19
+ specification does not include the ` [[sycl::device_has()]] ` attribute. The
20
+ published spec refers to this attribute as ` [[sycl::requires ()]] ` , but the
21
+ attribute was renamed to ` [[sycl::device_has()]] ` in [ this pull request ] [ 4 ] to
22
+ the SYCL 2020 specification .
25
23
26
24
[ 4 ] : < https://github.com/KhronosGroup/SYCL-Docs/pull/171 >
27
25
@@ -99,19 +97,19 @@ compiler must not issue any diagnostic simply because a kernel uses an optional
99
97
feature.
100
98
101
99
The only exception to this rule occurs when the application uses the C++
102
- attribute ` [[sycl::requires ()]] ` . When the application decorates a kernel or
100
+ attribute ` [[sycl::device_has ()]] ` . When the application decorates a kernel or
103
101
device function with this attribute, it is an assertion that the kernel or
104
102
device function is allowed to use only those optional features which are listed
105
103
by the attribute. Therefore, the FE compiler must issue a diagnostic if the
106
104
kernel or device function uses any other optional kernel features.
107
105
108
106
The SYCL 2020 specification only mandates this diagnostic when a kernel or
109
- device function that is decorated with ` [[sycl::requires ()]] ` uses an optional
110
- kernel feature (not listed in the attribute), ** and** when that use is in the
111
- kernel's static call graph as computed for the translation unit that contains
112
- the kernel function. Thus, the compiler is not required to issue a diagnostic
113
- if the use is in a ` SYCL_EXTERNAL ` function that is defined in another
114
- translation unit.
107
+ device function that is decorated with ` [[sycl::device_has ()]] ` uses an
108
+ optional kernel feature (not listed in the attribute), ** and** when that use is
109
+ in the kernel's static call graph as computed for the translation unit that
110
+ contains the kernel function. Thus, the compiler is not required to issue a
111
+ diagnostic if the use is in a ` SYCL_EXTERNAL ` function that is defined in
112
+ another translation unit.
115
113
116
114
Note that this behavior does not change when the compiler runs in AOT mode.
117
115
Even if the user specifies a target device via "-fsycl-targets", that does not
@@ -147,16 +145,17 @@ is required.
147
145
To further clarify, this exception must be thrown in the following
148
146
circumstances:
149
147
150
- * For a kernel that is not decorated with ` [[sycl::requires()]] ` , the exception
151
- must be thrown if the kernel uses a feature that the device does not support.
148
+ * For a kernel that is not decorated with ` [[sycl::device_has()]] ` , the
149
+ exception must be thrown if the kernel uses a feature that the device does
150
+ not support.
152
151
153
- * For a kernel that is decorated with ` [[sycl::requires ()]] ` , the exception
152
+ * For a kernel that is decorated with ` [[sycl::device_has ()]] ` , the exception
154
153
must be thrown if the device does not have the aspects listed in that
155
154
attribute. Note that the exception must be thrown even if the kernel does
156
155
not actually use a feature corresponding to the aspect, and it must be
157
156
thrown even if the aspect does not correspond to any optional feature.
158
157
159
- * For a kernel that is decorated with ` [[sycl::requires ()]] ` , the FE compiler
158
+ * For a kernel that is decorated with ` [[sycl::device_has ()]] ` , the FE compiler
160
159
will mostly check (at compile time) whether the kernel uses any features that
161
160
are not listed in the attribute. However, this check only results in a
162
161
warning, so the runtime is still responsible for throwing the exception if
@@ -198,8 +197,8 @@ of the required work-group or sub-group size).
198
197
199
198
As we will see later, it is helpful to decorate all APIs in DPC++ headers that
200
199
correspond to optional kernel features with a C++ attribute that identifies the
201
- associated aspect. We cannot use the ` [[sycl::requires ()]] ` attribute for this
202
- purpose, though, because that attribute is allowed only for functions.
200
+ associated aspect. We cannot use the ` [[sycl::device_has ()]] ` attribute for
201
+ this purpose, though, because that attribute is allowed only for functions.
203
202
Instead, we invent a new internal attribute ` [[sycl_detail::uses_aspects()]] `
204
203
that can be used to decorate both functions and types. This attribute is not
205
204
documented for user code; instead it is an internal implementation detail of
@@ -246,7 +245,7 @@ instantiations of `sycl::atomic_ref` as an optional feature.
246
245
247
246
Although the examples above show only a single aspect parameter to the
248
247
` [[sycl_detail::uses_aspects()]] ` attribute, this attribute should support a
249
- list of aspects, similar to the ` [[sycl::requires ()]] ` attribute. This will
248
+ list of aspects, similar to the ` [[sycl::device_has ()]] ` attribute. This will
250
249
allow us to support future features that depend on a conjunction of aspects
251
250
(e.g. a feature that does atomic operations on 64-bit floating point values
252
251
might be decorated with
@@ -261,7 +260,7 @@ attribute for any device code that uses the `double` type.
261
260
262
261
### New LLVM IR metadata
263
262
264
- In order to communicate the information from ` [[sycl::requires ()]] ` and
263
+ In order to communicate the information from ` [[sycl::device_has ()]] ` and
265
264
` [[sycl_detail::uses_aspects()]] ` attributes to the DPC++ post-link tool, we
266
265
introduce several new LLVM IR metadata.
267
266
@@ -295,18 +294,18 @@ type's name.
295
294
We also introduce two metadata that can be attached to a function definition
296
295
similar to the existing ` !intel_reqd_sub_group_size ` . The
297
296
` !intel_declared_aspects ` metadata is used for functions that are decorated
298
- with ` [[sycl::requires ()]] ` , and the ` !intel_used_aspects ` metadata is used to
299
- store the propagated information about all aspects used by a kernel or exported
300
- device function.
297
+ with ` [[sycl::device_has ()]] ` , and the ` !intel_used_aspects ` metadata is used
298
+ to store the propagated information about all aspects used by a kernel or
299
+ exported device function.
301
300
302
301
In each case, the metadata's parameter is an unnamed metadata node, and the
303
302
value of the metadata node is a list of ` i32 ` constants, where each constant is
304
303
a value from ` enum class aspect ` .
305
304
306
305
For example, the following illustrates the IR that corresponds to a function
307
- ` foo ` that is decorated with ` [[sycl::requires ()]] ` where the required aspects
308
- have the numerical values ` 8 ` and ` 9 ` . In addition, the function uses an
309
- optional feature that corresponds to an aspect with numerical value ` 8 ` .
306
+ ` foo ` that is decorated with ` [[sycl::device_has ()]] ` where the required
307
+ aspects have the numerical values ` 8 ` and ` 9 ` . In addition, the function uses
308
+ an optional feature that corresponds to an aspect with numerical value ` 8 ` .
310
309
311
310
```
312
311
define void @foo() !intel_declared_aspects !1 !intel_used_aspects !2 {}
@@ -318,7 +317,7 @@ define void @foo() !intel_declared_aspects !1 !intel_used_aspects !2 {}
318
317
### Changes to the DPC++ front-end
319
318
320
319
The front-end of the device compiler is responsible for parsing the
321
- ` [[sycl::requires ()]] ` and ` [[sycl_detail::uses_aspects()]] ` attributes and
320
+ ` [[sycl::device_has ()]] ` and ` [[sycl_detail::uses_aspects()]] ` attributes and
322
321
transferring the information to the LLVM IR metadata described above according
323
322
to the following rules:
324
323
@@ -331,7 +330,7 @@ to the following rules:
331
330
front-end adds an ` !intel_used_aspects ` metadata to the function's definition
332
331
listing the aspects from that attribute.
333
332
334
- * If a function is decorated with ` [[sycl::requires ()]] ` , the front-end adds
333
+ * If a function is decorated with ` [[sycl::device_has ()]] ` , the front-end adds
335
334
an ` !intel_declared_aspects ` metadata to the function's definition listing
336
335
the aspects from that attribute.
337
336
@@ -395,7 +394,7 @@ location information into the IR. To compensate, the warning message displays
395
394
the static call chain that leads to the problem. For example:
396
395
397
396
```
398
- warning: function 'foo' uses aspect 'fp64' not listed in 'sycl::requires '
397
+ warning: function 'foo' uses aspect 'fp64' not listed in 'sycl::device_has '
399
398
use is from this call chain:
400
399
foo()
401
400
bar()
@@ -412,7 +411,7 @@ When the compiler is invoked with `-g`, the implementation uses the
412
411
and column information like so:
413
412
414
413
```
415
- hw.cpp:27:4: warning: function 'foo' uses aspect 'fp64' not listed in 'sycl::requires '
414
+ hw.cpp:27:4: warning: function 'foo' uses aspect 'fp64' not listed in 'sycl::device_has '
416
415
use is from this call chain:
417
416
foo()
418
417
bar() hw.cpp:15:3
@@ -460,7 +459,7 @@ We do not propose this change as part of this design, though. We expect that
460
459
this will not be a common error because applications can avoid this problem by
461
460
declaring the ` SYCL_EXTERNAL ` function in a common header that is included by
462
461
both the importing and the exporting translation unit. If the declaration (in
463
- the header) is decorated with ` [[sycl::requires ()]] ` , the shared declaration
462
+ the header) is decorated with ` [[sycl::device_has ()]] ` , the shared declaration
464
463
will ensure that the definition stays in sync with the declaration.
465
464
466
465
@@ -492,7 +491,7 @@ For the purposes of this analysis, the set of *Used* aspects is computed by
492
491
taking the union of the aspects listed in the kernel's (or device function's)
493
492
` !intel_used_aspects ` and ` !intel_declared_aspects ` sets. This is consistent
494
493
with the SYCL specification, which says that a kernel decorated with
495
- ` [[sycl::requires ()]] ` may only be submitted to a device that provides the
494
+ ` [[sycl::device_has ()]] ` may only be submitted to a device that provides the
496
495
listed aspects, regardless of whether the kernel actually uses those aspects.
497
496
498
497
We must also split two kernels into different device images if they have
0 commit comments