@@ -244,7 +244,7 @@ In order to communicate the information from `[[sycl::requires()]]` and
244
244
introduce two new LLVM IR metadata that can be attached to a function
245
245
definition, similar to the existing ` !intel_reqd_sub_group_size ` .
246
246
247
- These new metadata are named ` !intel_allowed_aspects ` and
247
+ These new metadata are named ` !intel_declared_aspects ` and
248
248
` !intel_used_aspects ` . In each case, the parameter is an (unnamed) metadata
249
249
node, and the value of the metadata node is a list of ` i32 ` constants, where
250
250
each constant is a value from ` enum class aspect ` . For example, the following
@@ -254,7 +254,7 @@ illustrates the IR that corresponds to a function `foo` that is decorated with
254
254
corresponds to an aspect with numerical value ` 8 ` .
255
255
256
256
```
257
- define void @foo() !intel_allowed_aspects !1 !intel_used_aspects !2 {}
257
+ define void @foo() !intel_declared_aspects !1 !intel_used_aspects !2 {}
258
258
!1 = !{i32 8, i32 9}
259
259
!2 = !{i32 8}
260
260
```
@@ -264,11 +264,11 @@ define void @foo() !intel_allowed_aspects !1 !intel_used_aspects !2 {}
264
264
265
265
The front-end of the device compiler is responsible for parsing the
266
266
` [[sycl::requires()]] ` and ` [[sycl_detail::uses_aspects()]] ` attributes and
267
- transferring the information to the LLVM IR ` !intel_allowed_aspects ` and
267
+ transferring the information to the LLVM IR ` !intel_declared_aspects ` and
268
268
` !intel_used_aspects ` metadata according to the following rules:
269
269
270
270
* If a function is decorated with the ` [[sycl::requires()]] ` attribute, the
271
- front-end emits an ` !intel_allowed_aspects ` metadata on the function's LLVM
271
+ front-end emits an ` !intel_declared_aspects ` metadata on the function's LLVM
272
272
IR definition with the numerical values of the aspects listed in the
273
273
attribute.
274
274
@@ -332,7 +332,7 @@ already know whether an expression is potentially evaluated.
332
332
### Changes to other phases of clang
333
333
334
334
Any clang phases that do function inlining will need to be changed, so that the
335
- ` !intel_allowed_aspects ` and ` !intel_uses_aspects ` metadata are transferred
335
+ ` !intel_declared_aspects ` and ` !intel_uses_aspects ` metadata are transferred
336
336
from the inlined function to the function that receives the inlined function
337
337
body. Presumably, there is already similar logic for the existing
338
338
` !reqd_work_group_size ` metadata, which already decorates device functions.
@@ -376,24 +376,24 @@ Linking][5].
376
376
377
377
This pass operates on the static call graph for each kernel and each exported
378
378
device function, propagating the aspects from the ` !intel_used_aspects ` and
379
- ` !intel_allowed_aspects ` metadata from the leaves of the call graph up to their
380
- callers. The result of this pass is that each device function is labeled with
381
- a * Used* set of aspects which is computed as the union of the following:
379
+ ` !intel_declared_aspects ` metadata from the leaves of the call graph up to
380
+ their callers. The result of this pass is that each device function is labeled
381
+ with a * Used* set of aspects which is computed as the union of the following:
382
382
383
383
* The aspects in the function's ` !intel_used_aspects ` metadata (if any).
384
- * The aspects in the function's ` !intel_allowed_aspects ` metadata (if any).
384
+ * The aspects in the function's ` !intel_declared_aspects ` metadata (if any).
385
385
* The aspects in the * Used* set of all functions called by this function.
386
386
387
387
Once the * Used* set of aspects is known for each function, the post-link tool
388
- compares this set of aspects with the aspects from any ` !intel_allowed_aspects `
389
- metadata. If the function has this metadata and if the * Used * set contains
390
- aspects not in that set, it issues a warning indicating that the function uses
391
- aspects that are not in the ` [[sycl::requires()]] ` list. Unfortunately, the
392
- post-link tool is unable to include the source position of the code that uses
393
- the aspect in question. To compensate, the warning message must include
394
- instructions telling the user how to run the clang static analyzer which
395
- provides a better diagnostic. This analysis phase is described in more detail
396
- below.
388
+ compares this set of aspects with the aspects from any
389
+ ` !intel_declared_aspects ` metadata. If the function has this metadata and if
390
+ the * Used * set contains aspects not in that set, it issues a warning indicating
391
+ that the function uses aspects that are not in the ` [[sycl::requires()]] ` list.
392
+ Unfortunately, the post-link tool is unable to include the source position of
393
+ the code that uses the aspect in question. To compensate, the warning message
394
+ must include instructions telling the user how to run the clang static analyzer
395
+ which provides a better diagnostic. This analysis phase is described in more
396
+ detail below.
397
397
398
398
#### Changes to the device code split algorithm
399
399
0 commit comments