Skip to content

Commit 4ec2881

Browse files
committed
Address more review comments
* Capture some more open issues with properties that are represented as IR attributes on kernel arguments. What happens if an single aggregate kernel argument gets properties from more than one source? * Resolve TODO about replacing `handler::kernel_single_task()` with a wrapper class like `KernelSingleTaskWrapper`. The front-end team thinks this is the preferred direction. * Add note that `<initializer_list>` must be included in order to use the optional "filter list" parameter to the C++ attributes. This "filter list" parameter is a brace-enclosed list, and the front-end team thinks it would be easier to implement if `<initializer_list>` is included.
1 parent 4476369 commit 4ec2881

File tree

1 file changed

+38
-16
lines changed

1 file changed

+38
-16
lines changed

sycl/doc/CompileTimeProperties.md

Lines changed: 38 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -89,12 +89,13 @@ class
8989
```
9090

9191
The `[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute has an
92-
even number of parameters. The first half of the parameters are the names of
93-
the properties, and the second half of the parameters are the values for those
94-
properties. Each property has exactly one value, so the property at parameter
95-
position 0 corresponds to the value at position _N / 2_, etc. To illustrate
96-
using the same example as before, the result of the parameter pack expansion
97-
would look like this:
92+
even number of parameters, assuming that the optional "filter list" parameter
93+
is not specified (see below for a description of this parameter). The first
94+
half of the parameters are the names of the properties, and the second half of
95+
the parameters are the values for those properties. Each property has exactly
96+
one value, so the property at parameter position 0 corresponds to the value at
97+
position _N / 2_, etc. To illustrate using the same example as before, the
98+
result of the parameter pack expansion would look like this:
9899

99100
```
100101
namespace sycl::ext::oneapi {
@@ -244,12 +245,28 @@ the property value to a string if it is not already a string.
244245

245246
[7]: <https://llvm.org/doxygen/classllvm_1_1Function.html#a092beb46ecce99e6b39628ee92ccd95a>
246247

247-
**TODO**: What happens when a "sycl special class" object is captured as a
248-
kernel argument? The compiler passes each member of the class as a separate
249-
argument. Should the device compiler duplicate the properties on each such
250-
parameter in this case? Or, is it the header's responsibility to add the C++
251-
attribute to one of the member variables in this case? How does the header
252-
decide which member variable to decorate, though?
248+
**TODO**: There are a number of open issues with this attribute and with the
249+
semantics of properties that are represented as attributes on kernel
250+
arguments. Suppose there are two SYCL types that take properties: _A_ and
251+
_B_. (For example, this could be two specializations of `annotated_ptr`, each
252+
decorated with different properties.) Now suppose the application creates a
253+
struct that contains members with both of these types, and it passes that
254+
struct as a kernel argument. What is the intended semantic? Does the argument
255+
get decorated with the union of the properties on both _A_ and _B_? What if
256+
those properties are mutually exclusive? A similar case exists when the
257+
application creates a struct that inherits from both _A_ and _B_.
258+
259+
The previous example shows a case when a single kernel argument gets properties
260+
from two (or more) types. However, the opposite can also occur. Certain SYCL
261+
classes are decorated with `__attribute__((sycl_special_class))`, which causes
262+
the compiler to pass each member of that class as a separate kernel argument.
263+
What should happen with the properties that decorate the class? Should the
264+
compiler duplicate the properties on each such kernel argument? Or, maybe it
265+
should be the header file's responsibility not to decorate such a class with
266+
`[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]`, and instead it
267+
should decorate specific member variable(s) with this attribute? How does the
268+
header decide which properties are used to decorate which member variables,
269+
though?
253270

254271

255272
## Properties on kernel functions
@@ -344,11 +361,12 @@ string if it is not already a string.
344361

345362
[9]: <https://llvm.org/doxygen/classllvm_1_1Function.html#ae7b919df259dce5480774e656791c079>
346363

347-
**TODO**: The intention is to replace the existing member functions like
364+
**NOTE**: The intention is to replace the existing member functions like
348365
`handler::kernel_single_task()` with wrapper classes like
349-
`KernelSingleTaskWrapper`. Does this pose any problems? There are comments in
350-
the headers indicating that the front-end recognizes the function
351-
`handler::kernel_single_task()` by name.
366+
`KernelSingleTaskWrapper`. We believe this will not cause problems for the
367+
device compiler front-end because it recognizes kernel functions via the
368+
`__attribute__((sycl_kernel))` attribute, not by the name
369+
`handler::kernel_single_task()`.
352370

353371

354372
## Properties on a non-global variable type
@@ -618,6 +636,10 @@ first parameter that is a brace-enclosed list of property names:
618636
* `[[__sycl_detail__::add_ir_function_attributes()]]`
619637
* `[[__sycl_detail__::add_ir_member_annotation()]]`
620638

639+
Since this brace-enclosed list acts somewhat like an initializer list, the
640+
header must include `<initializer_list>` prior to passing this optional first
641+
parameter.
642+
621643
The front-end treats this list as a "pass list", ignoring any property whose
622644
name is not in the list. To illustrate, consider the following example where
623645
`accessor` treats some properties as "kernel parameter attributes" and others

0 commit comments

Comments
 (0)