From ba0d71f2c3b3af1f717b4839103933ae095b6f72 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Tue, 26 Oct 2021 17:43:51 -0400 Subject: [PATCH 1/7] [SYCL][Doc] Design doc for compile-time properties Add a design document for the DPC++ extension API `SYCL_EXT_ONEAPI_PROPERTY_LIST`, which describes how compile-time properties are recognized by the front-end, how they are represented in LLVM IR, and how they are translated into SPIR-V. --- sycl/doc/CompileTimeProperties.md | 680 ++++++++++++++++++++++++++++++ 1 file changed, 680 insertions(+) create mode 100644 sycl/doc/CompileTimeProperties.md diff --git a/sycl/doc/CompileTimeProperties.md b/sycl/doc/CompileTimeProperties.md new file mode 100644 index 0000000000000..27500e808c5c7 --- /dev/null +++ b/sycl/doc/CompileTimeProperties.md @@ -0,0 +1,680 @@ +# Implementation design for compile time constant properties + +This document describes the implementation design for the DPC++ extension +[SYCL\_EXT\_ONEAPI\_PROPERTY\_LIST][1], which adds a general mechanism for +specifying properties which are known at compile time. This extension is not +itself a feature, but rather a building block that can be incorporated into +other features. + +[1]: + +There are a number of situations where we plan to use compile-time constant +properties, but this design document does not attempt to address them all. +Rather, it describes the design for each "category" of use and illustrates each +category with a specific feature. For example `accessor` is used to illustrate +properties that are applied to a kernel argument, but the same technique could +be used for other variables that are captured as kernel arguments. + +In all cases, the goal of this design is to explain how a DPC++ program that +uses properties is consumed by the device compiler and eventually represented +in LLVM IR. This typically involves some logic in the header files which +results in a C++ annotation that contains the properties. The device compiler +front-end is responsible for consuming this annotation and producing some +corresponding LLVM IR. One of the goals of this design is to avoid changes to +the front-end each time we add a new property, so the front-end is not required +to understand each property it consumes. Instead, it follows a mechanical +process for converting properties listed in the C++ annotation into LLVM IR, +and this mechanical process need not be updated when we add new properties. + +Once the information about properties is represented in IR, it is available to +compiler passes. For example, the `sycl-post-link` tool might use a property +in order to perform one of its transformations. Some properties are consumed +by the DPC++ compiler, but others are transformed into SPIR-V for use by the +JIT compiler. This design document also describes how this SPIR-V +transformation is done. + + +## Properties on a global variable type + +One use for compile-time properties is with types that are used exclusively +for declaring global variables. One such example is the +[SYCL\_EXT\_ONEAPI\_DEVICE\_GLOBAL][2] extension: + +[2]: + +``` +namespace sycl::ext::oneapi { + +template > +class device_global {/*...*/}; + +} // namespace sycl::ext::oneapi +``` + +The following code illustrates a `device_global` variable that is declared with +two compile-time properties: + +``` +using sycl::ext::oneapi; + +device_global>> + dm1; +``` + +The header file represents these properties with an internal C++ attribute +named `[[__sycl_detail__::add_ir_global_variable_attributes()]]` whose value +is a list that is created through a template parameter pack expansion: + +``` +namespace sycl::ext::oneapi { + +template > +class device_global {/*...*/}; + +// Partial specialization to make PropertyListT visible as a parameter pack +// of properties. +template +class +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_global_variable_attributes( + Props::meta_name..., Props::meta_value... + )]] +#endif + device_global> {/*...*/}; + +} // namespace sycl::ext::oneapi +``` + +The initial entries in the C++ attribute's parameter list are the names of the +properties, and these are followed by the values of the properties. To +illustrate using the same example as before, the result of the parameter pack +expansion would look like this: + +``` +namespace sycl::ext::oneapi { + +template class +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_global_variable_attributes( + "sycl-device-image-scope", // Name of first property + "sycl-host-access", // Name of second property + "", // First property has no parameter + "read" // Value of second property + )]] +#endif + device_global {/*...*/}; + +} // namespace sycl::ext::oneapi +``` + +The device compiler only uses the +`[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute when the +decorated type is used to create an [LLVM IR global variable][3] and the global +variable's type is either: + +* The type that is decorated by the attribute, or +* An array of the type that is decorated by the attribute. + +[3]: + +The device compiler front-end silently ignores the attribute when the decorated +type is used in any other way. + +When the device compiler front-end creates a global variable from the decorated +type as described above, it also adds one IR attribute to the global variable +for each property using +[`GlobalVariable::addAttribute(StringRef, StringRef)`][4]. If the property +value is not already a string, it converts it to a string as described in +[Property representation in C++ attributes][5]. + +[4]: +[5]: <#property-representation-in-C-attributes> + +Note that the front-end does not need to understand any of the properties in +order to do this translation. + + +## Properties on kernel arguments + +Another use of compile-time properties is with types that are used to define +kernel arguments. For example, the [SYCL\_ONEAPI\_accessor\_properties][6] +extension could be redesigned to use compile-time properties. Such a redesign +might look like: + +[6]: + +``` +namespace sycl { + +template > +class accessor {/* ... */}; + +} // namespace sycl +``` + +Typical usage would look like this (showing a hypothetical property named +`foo`): + +``` +using sycl; +using sycl::ext::oneapi; + +accessor acc(buf, cgh, property_list{no_alias_v, foo_v<32>}); +``` + +As before, the header file represents the properties with an internal C++ +attribute, where the initial parameters are the names of the properties and +the subsequent parameters are the property values. + +``` +namespace sycl { + +template > +class accessor {/* ... */}; + +// Partial specialization to make PropertyListT visible as a parameter pack +// of properties. +template +class +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_kernel_parameter_attributes( + Props::meta_name..., Props::meta_value... + )]] +#endif + accessor> {/*...*/}; + +} // namespace sycl +``` + +Illustrating this with the previous example: + +``` +namespace sycl { + +template class +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_kernel_parameter_attributes( + "sycl-no-alias", // Name of first property + "sycl-foo", // Name of second property + "", // First property has no parameter + 32 // Value of second property + )]] +#endif + accessor {/* ... */}; + +} // namespace sycl +``` + +As the name of the C++ attribute suggests, the device compiler front-end uses +the attribute only when the decorated type is the type of a kernel argument, +and it silently ignores the attribute when the class is used in any other way. + +When the device compiler front-end creates a kernel argument in this way, it +adds one LLVM IR attribute to the kernel function's parameter for each property +in the list. For example, this can be done by calling +[`Function::addParamAttrs(unsigned ArgNo, const AttrBuilder &)`][7]. As +before, the IR attributes are added as strings, so the front-end must convert +the property value to a string if it is not already a string. + +[7]: + +**TODO**: What happens when a "sycl special class" object is captured as a +kernel argument? The compiler passes each member of the class as a separate +argument. Should the device compiler duplicate the properties on each such +parameter in this case? Or, is it the header's responsibility to add the C++ +attribute to one of the member variables in this case? How does the header +decide which member variable to decorate, though? + + +## Properties on kernel functions + +Compile-time properties can also be used to decorate kernel functions as with +the [SYCL\_EXT\_ONEAPI\_KERNEL\_PROPERTIES][8] extension. There are two ways +the application can specify these properties. The first is by passing a +`property_list` parameter to the function that submits the kernel: + +[8]: + +``` +namespace sycl { + +class handler { + template + void single_task(PropertyListT properties, const KernelType &kernelFunc); +}; + +// namespace sycl +``` + +For example: + +``` +using sycl; +using sycl::ext::oneapi; + +void foo(handler &cgh) { + cgh.single_task( + property_list{sub_group_size_v<32>, device_has_v}, + [=] {/* ... */}); +} +``` + +The second way an application can specify kernel properties is by adding a +`properties` member variable to a named kernel function object: + +``` +using sycl; +using sycl::ext::oneapi; + +class MyKernel { + public: + void operator()() {/* ... */} + + static constexpr auto properties = + property_list{sub_group_size_v<32>, device_has_v}; +}; + +void foo(handler &cgh) { + MyKernel k; + cgh.single_task(k); +} +``` + +Internally, the headers lower both cases to a wrapper class that expresses the +properties as an internal C++ attribute, and the `operator()` of this class +becomes the "top level" kernel function that is recognized by the front-end. + +``` +template +class KernelSingleTaskWrapper; + +// Partial specialization to make PropertyListT visible as a parameter pack +// of properties. +template +class KernelSingleTaskWrapper> { + KernelType k; + + public: + KernelSingleTaskWrapper(KernelType k) : k(k) {} + +#ifdef __SYCL_DEVICE_ONLY__ + __attribute__((sycl_kernel)) + [[__sycl_detail__::add_ir_function_attributes( + Props::meta_name..., Props::meta_value... + )]] +#endif + void operator()() {k();} +}; +``` + +Although the DPC++ headers only use the +`[[__sycl_detail__::add_ir_function_attributes()]]` attribute on the definition +of a kernel function as shown above, the front-end recognizes it for any +function definition. The front-end adds one LLVM IR function attribute for +each property in the list. For example, this can be done by calling +[`Function::addFnAttr(StringRef, StringRef)`][9]. As before, the IR attributes +are added as strings, so the front-end must convert the property value to a +string if it is not already a string. + +[9]: + +**TODO**: The intention is to replace the existing member functions like +`handler::kernel_single_task()` with wrapper classes like +`KernelSingleTaskWrapper`. Does this pose any problems? There are comments in +the headers indicating that the front-end recognizes the function +`handler::kernel_single_task()` by name. + + +## Properties on a non-global variable type + +Another use of compile-time properties is with types that are used to define +non-global variables. An example of this is the proposed `annotated_ptr` +class. + +``` +namespace sycl::ext::oneapi { + +template > +class annotated_ptr { + T *ptr; + public: + annotated_ptr(T *p) : ptr(p) {} +}; + +} // namespace sycl::ext::oneapi +``` + +where an example use looks like: + +``` +using sycl::ext::oneapi; + +void foo(int *p) { + annotated_ptr>> + aptr(p); +} +``` + +We again implement the property list in the header via a C++ attribute, though +this time the attribute decorates a member variable of the class: + +``` +namespace sycl::ext::oneapi { + +template > +class annotated_ptr; + +// Partial specialization to make PropertyListT visible as a parameter pack +// of properties. +template +class annotated_ptr> { + T *ptr +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_member_annotation( + Props::meta_name..., Props::meta_value... + )]] +#endif + ; + public: + annotated_ptr(T *p) : ptr(p) {} +}; + +} // namespace sycl::ext::oneapi +``` + +Illustrating this with properties from our previous example: + +``` +namespace sycl::ext::oneapi { + +template > +class annotated_ptr; + +// Partial specialization to make PropertyListT visible as a parameter pack +// of properties. +template +class annotated_ptr> { + T *ptr +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_member_annotation( + "foo", // Name of first property + "bar", // Name of second property + "", // First property has no parameter + 32 // Value of second property + )]] +#endif + ; + public: + annotated_ptr(T *p) : ptr(p) {} +}; + +} // namespace sycl::ext::oneapi +``` + +When the device compiler generates code to reference the decorated member +variable, it emits a call to the LLVM intrinsic function +[`@llvm.ptr.annotation`][10] that annotates the pointer to that member +variables, similar to the way the existing clang `__attribute__((annotate()))` +works. Illustrating this with some simplified LLVM IR that matches the example +code above: + +[10]: + +``` +@.str = private unnamed_addr constant [27 x i8] c"sycl-properties:foo,bar=32\00", + section "llvm.metadata" +@.str.1 = private unnamed_addr constant [9 x i8] c"file.cpp\00", + section "llvm.metadata" + +define void @foo(i32* %ptr) { + %aptr = alloca %class.annotated_ptr + %ptr = getelementptr inbounds %class.annotated_ptr, %class.annotated_ptr* %aptr, + i32 0, i32 0 + %1 = bitcast i32** %ptr to i8* + %2 = call i8* @llvm.ptr.annotation.p0i8(i8* %1, + i8* getelementptr inbounds ([27 x i8], [27 x i8]* @.str, i32 0, i32 0), + i8* getelementptr inbounds ([9 x i8], [9 x i8]* @.str.1, i32 0, i32 0), + i32 3, i8* null) + %3 = bitcast i8* %2 to i32** + store i32* %ptr, i32** %3 + ret void +} +``` + +The front-end encodes the properties from the C++ attribute +`[[__sycl_detail__::add_ir_member_annotation()]]` into the annotation string +(`@.str` in the example above) using the following algorithm: + +* The property value is converted to a string as specified in + [Property representation in C++ attributes][5]. +* Construct a property definition string for each property: + - If the property value is the empty string, the property definition is just + the name of the property. + - Otherwise, the property definition string is formed by concatenating the + property name with the equal sign (`=`) and the property value. +* The annotation string is formed by concatenating all property definition + strings, separated by a comma (`,`). +* The annotation string is pre-pended with `"sycl-properties:"` and NULL + terminated. + +**NOTE**: Calls to the `@llvm.ptr.annotation` intrinsic function are known to +disable many clang optimizations. As a result, properties added to a +non-global variable will likely result in LLVM IR (and SPIR-V) that is not well +optimized. This puts more pressure on the SPIR-V consumer (e.g. JIT compiler) +to perform these optimizations. + + +## Property representation in C++ attributes + +As noted above, there are several C++ attributes that convey property names and +values to the front-end: + +* `[[__sycl_detail__::add_ir_global_variable_attributes()]]` +* `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` +* `[[__sycl_detail__::add_ir_function_attributes()]]` +* `[[__sycl_detail__::add_ir_member_annotation()]]` + +All of these attributes take a parameter list with the same format. There are +always an even number of parameters, where the first half are the property +names and the second half are the property values. The property name is always +a string literal or a `constexpr char *` expression. By convention, property +names that correspond to LLVM IR attributes normally start with the prefix +`"sycl-"` in order to avoid collision with non-SYCL IR attributes, but this is +not a strict requirement. + +The property value can be a literal or `constexpr` expression of the following +types: + +* `const char *`. +* An integer type. +* A floating point type. +* A boolean type. +* A character type. +* An enumeration type. + +All properties require a value when represented in the C++ attribute. If the +SYCL property has no value the header passes the empty string (`""`). + +The front-end converts each value to a string before representing it in LLVM +IR. Integer and floating point values are converted with the same format as +`std::to_string()` would produce. Boolean values are converted to either +`"true"` or `"false"`. Enumeration values are first converted to an integer +and then converted to a string with the same format as `std::to_string()`. + +**TODO**: Should we allow property values that are type names? If so, I +suppose they would be converted to a string representation of the mangled name? + +**TODO**: Should we allow property values of other (non-fundamental) types? If +we allow this, we need to teach the front-end how to convert each type to a +string, which means the front-end needs to be changed each time we add a +property with a new non-fundamental type. This seems undesirable. However, if +we do not allow non-fundamental types, how do we represent properties like +`work_group_size`, whose value is a 3-tuple of integers? Maybe we could just +allow `std::tuple`, where the type of each element is one of the fundamental +types listed above. + + +## Representing properties in SPIR-V + +There is no mechanical process which converts all LLVM IR attributes to +SPIR-V. This is because we do not need all properties to be expressed in +SPIR-V and because there is no consistent way to represent properties in +SPIR-V. Therefore, the `sycl-post-link` tool decides on a case-by-case basis +which properties are translated into SPIR-V and which representation to use. + +We use the [SPIR-V LLVM Translator][11] to translate from LLVM IR to SPIR-V, +and that tool defines [idiomatic LLVM IR][12] representations that correspond +to various SPIR-V instructions. Therefore, the `sycl-post-link` tool can +translate a property into a specific SPIR-V instruction by generating the +corresponding idiomatic LLVM IR. The following sections describe some common +cases. + +[11]: +[12]: + +### Property on a kernel function + +When a property on a kernel function needs to be represented in SPIR-V, we +generally translate the property into a SPIR-V **OpExecutionMode** instruction. +The SPIR-V LLVM Translator has an existing way to generate this instruction +when the LLVM IR contains the named metadata `!spirv.ExecutionMode` as +illustrated below: + +``` +!spirv.ExecutionMode = !{!0, !1} ; Each operand in this metadata + ; represents one OpExectionMode + ; instruction that will be generated. +!0 = !{void ()* @bar, i32 42} ; The first operand identifies a kernel + ; function. The second operand is the + ; integer value of a SPIR-V execution + ; mode. +!1 = !{void ()* @bar, i32 43, i32 3} ; Any additional operands in the metadata + ; correspond to "extra operands" to the + ; OpExecutionMode instruction. These + ; operands must be integer literals. +``` + +### Property on a kernel parameter + +When a property on a kernel parameter needs to be represented in SPIR-V, we +generally translate the property into a SPIR-V **OpDecorate** instruction for +the corresponding **OpFunctionParameter** of the kernel function. Since the +SPIR-V LLVM Translator does not have an existing way to generate these +decorations, we propose the following mechanism. + +An LLVM IR function definition may optionally have a metadata kind of +`!spirv.ParameterDecorations`. If it does, that metadata node must have one +operand for each of the function's parameters. Each of those operands is +another metadata node that describes the decorations for that parameter. To +illustrate: + +``` +define spir_kernel void @MyKernel(%arg1, %arg2) !spirv.ParameterDecorations !0 { +} + +!0 = !{!1, !2} ; Each operand in this metadata represents the + ; decorations for one kernel parameter. +!1 = !{!3, !4} ; The first kernel parameter has two decorations. +!2 = !{} ; The second kernel parameter has no decorations. +!3 = !{i32 7742} ; This is the integer value of the first decoration. +!4 = !{i32 7743, i32 10} ; The first operand is the integer value of the + ; second decoration. Additional operands are + ; "extra operands" to the decoration. These + ; operands may be either integer literals or string + ; literals. +``` + +### Property on a global variable + +When a property on a global variable needs to be represented in SPIR-V, we +generally translate the property into a SPIR-V **OpDecorate** instruction for +the corresponding module scope (global) **OpVariable**. Again, there is no +existing mechanism to do this in the SPIR-V LLVM Translator, so we propose the +following mechanism. + +An LLVM IR global variable definition may optionally have a metadata kind of +`!spirv.Decorations`. If it does, that metadata node has one operand for each +of the global variable's decorations. To illustrate: + +``` +@MyVariable = global %MyClass !spirv.Decorations !0 +!0 = !{!1, !2} ; Each operand in this metadata represents one + ; decoration on the variable. +!1 = !{i32 7744} ; This is the integer value of the first decoration. +!2 = !{i32 7745, i32 20} ; The first operand is the integer value of the + ; second decoration. Additional operands are + ; "extra operands" to the decoration. These + ; operands may be either integer literals or string + ; literals. +``` + +### Property on a structure member of a non-global variable + +As we noted earlier, a property on a structure member variable is represented +in LLVM IR as a call to the intrinsic function `@llvm.ptr.annotation`, where +the annotation string starts with the prefix `"sycl-properties:"`. In order to +understand how these SYCL properties are translated into SPIR-V, it's useful to +review how a normal (i.e. non-SYCL) call to `@llvm.ptr.annotation` is +translated. + +The existing behavior of the SPIR-V LLVM Translator is to translate this call +into one (or both) of the following: + +* An **OpDecorate** instruction that decorates the intermediate pointer value + that is returned by the intrinsic (i.e. the pointer to the member variable). + +* An **OpMemberDecorate** instruction that decorates the member variable + itself. + +In both cases, the decoration is a single **UserSemantic** decoration where the +string literal is the same as the string literal in the LLVM annotation. + +When a SYCL structure member property needs to be represented in SPIR-V, +however, we prefer to represent each property as an extended SPIR-V decoration +rather than using a **UserSemantic** decoration. There is no existing +mechanism in the SPIR-V LLVM Translator to generate extended decorations like +this, so we propose the following new mechanism. + +When a member variable property needs to be represented in SPIR-V, the +`sycl-post-link` tool converts the `@llvm.ptr.annotation` intrinsic call into a +call to the SPIR-V intrinsic `__spirv_AddMemberDecoration` which has a metadata +function argument that specifies the decorations as illustrated below: + +``` +%annotated_ptr = call i8* __spirv_AddMemberDecoration(i8* %ptr, metadata !0) + +!0 = !{!1, !2} ; Each operand in this metadata represents one + ; decoration. +!1 = !{i32 7744} ; This is the integer value of the first decoration. +!2 = !{i32 7745, i32 20} ; The first operand is the integer value of the + ; second decoration. Additional operands are + ; "extra operands" to the decoration. These + ; operands may be either integer literals or string + ; literals. +``` From 44763693264bc91488fc5c8772b61f0cc5d9210b Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Wed, 8 Dec 2021 13:32:02 -0500 Subject: [PATCH 2/7] Address review comments * When representing properties using `@llvm.ptr.annotation`, represent all properties in the fifth argument. This allows each property and its value to be represented as its own metadata, rather than combining them all into a single string. * Add an initial optional parameter to each C++ attribute that allows filtering of the properties. * The header now passes `nullptr` instead of `""` to represent the "value" of a property that has no value. * Clarify that each property in the C++ attribute parameter list has exactly one value, so the number of parameters is even (assuming the initial optional parameter is not specified). --- sycl/doc/CompileTimeProperties.md | 216 +++++++++++++++++++++++------- 1 file changed, 169 insertions(+), 47 deletions(-) diff --git a/sycl/doc/CompileTimeProperties.md b/sycl/doc/CompileTimeProperties.md index 27500e808c5c7..8b11f4dad9f0b 100644 --- a/sycl/doc/CompileTimeProperties.md +++ b/sycl/doc/CompileTimeProperties.md @@ -88,10 +88,13 @@ class } // namespace sycl::ext::oneapi ``` -The initial entries in the C++ attribute's parameter list are the names of the -properties, and these are followed by the values of the properties. To -illustrate using the same example as before, the result of the parameter pack -expansion would look like this: +The `[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute has an +even number of parameters. The first half of the parameters are the names of +the properties, and the second half of the parameters are the values for those +properties. Each property has exactly one value, so the property at parameter +position 0 corresponds to the value at position _N / 2_, etc. To illustrate +using the same example as before, the result of the parameter pack expansion +would look like this: ``` namespace sycl::ext::oneapi { @@ -101,7 +104,7 @@ template class [[__sycl_detail__::add_ir_global_variable_attributes( "sycl-device-image-scope", // Name of first property "sycl-host-access", // Name of second property - "", // First property has no parameter + nullptr, // First property has no parameter "read" // Value of second property )]] #endif @@ -128,10 +131,10 @@ type as described above, it also adds one IR attribute to the global variable for each property using [`GlobalVariable::addAttribute(StringRef, StringRef)`][4]. If the property value is not already a string, it converts it to a string as described in -[Property representation in C++ attributes][5]. +[IR representation as IR attributes][5]. [4]: -[5]: <#property-representation-in-C-attributes> +[5]: <#ir-representation-as-ir-attributes> Note that the front-end does not need to understand any of the properties in order to do this translation. @@ -219,7 +222,7 @@ template class [[__sycl_detail__::add_ir_kernel_parameter_attributes( "sycl-no-alias", // Name of first property "sycl-foo", // Name of second property - "", // First property has no parameter + nullptr, // First property has no parameter 32 // Value of second property )]] #endif @@ -423,10 +426,10 @@ class annotated_ptr> { T *ptr #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_member_annotation( - "foo", // Name of first property - "bar", // Name of second property - "", // First property has no parameter - 32 // Value of second property + "sycl-foo", // Name of first property + "sycl-bar", // Name of second property + nullptr, // First property has no parameter + 32 // Value of second property )]] #endif ; @@ -440,27 +443,41 @@ class annotated_ptr> { When the device compiler generates code to reference the decorated member variable, it emits a call to the LLVM intrinsic function [`@llvm.ptr.annotation`][10] that annotates the pointer to that member -variables, similar to the way the existing clang `__attribute__((annotate()))` +variables, similar to the way the existing `[[clang::annotate()]]` attribute works. Illustrating this with some simplified LLVM IR that matches the example code above: [10]: ``` -@.str = private unnamed_addr constant [27 x i8] c"sycl-properties:foo,bar=32\00", - section "llvm.metadata" +@.str = private unnamed_addr constant [16 x i8] c"sycl-properties\00", + section "llvm.metadata" @.str.1 = private unnamed_addr constant [9 x i8] c"file.cpp\00", - section "llvm.metadata" + section "llvm.metadata" +@.str.2 = private unnamed_addr constant [9 x i8] c"sycl-foo\00", align 1 +@.str.3 = private unnamed_addr constant [9 x i8] c"sycl-bar\00", align 1 + +@.args = private unnamed_addr constant { [9 x i8]*, i8*, [9 x i8]*, i32 } + { + [9 x i8]* @.str.2, ; Name of first property "sycl-foo" + i8* null, ; Null indicates this property has no value + [9 x i8]* @.str.3, ; Name of second property "sycl-bar" + i32 32 ; Value of second property + }, + section "llvm.metadata" define void @foo(i32* %ptr) { %aptr = alloca %class.annotated_ptr %ptr = getelementptr inbounds %class.annotated_ptr, %class.annotated_ptr* %aptr, i32 0, i32 0 %1 = bitcast i32** %ptr to i8* - %2 = call i8* @llvm.ptr.annotation.p0i8(i8* %1, - i8* getelementptr inbounds ([27 x i8], [27 x i8]* @.str, i32 0, i32 0), - i8* getelementptr inbounds ([9 x i8], [9 x i8]* @.str.1, i32 0, i32 0), - i32 3, i8* null) + + %2 = call i8* @llvm.ptr.annotation.p0i8(i8* nonnull %0, + i8* getelementptr inbounds ([16 x i8], [16 x i8]* @.str, i64 0, i64 0), + i8* getelementptr inbounds ([9 x i8], [9 x i8]* @.str.1, i64 0, i64 0), + i32 3, + i8* bitcast ({ [9 x i8]*, i8*, [9 x i8]*, i32 }* @.args to i8*)) + %3 = bitcast i8* %2 to i32** store i32* %ptr, i32** %3 ret void @@ -468,20 +485,21 @@ define void @foo(i32* %ptr) { ``` The front-end encodes the properties from the C++ attribute -`[[__sycl_detail__::add_ir_member_annotation()]]` into the annotation string -(`@.str` in the example above) using the following algorithm: - -* The property value is converted to a string as specified in - [Property representation in C++ attributes][5]. -* Construct a property definition string for each property: - - If the property value is the empty string, the property definition is just - the name of the property. - - Otherwise, the property definition string is formed by concatenating the - property name with the equal sign (`=`) and the property value. -* The annotation string is formed by concatenating all property definition - strings, separated by a comma (`,`). -* The annotation string is pre-pended with `"sycl-properties:"` and NULL - terminated. +`[[__sycl_detail__::add_ir_member_annotation()]]` into the +`@llvm.ptr.annotation` call as follows: + +* The first parameter to `@llvm.ptr.annotation` is the pointer to annotate (as + with any call to this intrinsic). +* The second parameter is the literal string `"sycl-properties"`. +* The third parameter is the name of the source file (as with any call to this + intrinsic). +* The fourth parameter is the line number (as with any call to this intrinsic). +* The fifth parameter is a metadata tuple with information about all of the + properties. The first element of the tuple is a string literal with the name + of the first property. The second element is the value of the first + property. The third element is a string literal with the name of the second + property, etc. Since each property has exactly one value, this tuple has an + even number of elements. **NOTE**: Calls to the `@llvm.ptr.annotation` intrinsic function are known to disable many clang optimizations. As a result, properties added to a @@ -490,7 +508,7 @@ optimized. This puts more pressure on the SPIR-V consumer (e.g. JIT compiler) to perform these optimizations. -## Property representation in C++ attributes +## Property representation in C++ attributes and in IR As noted above, there are several C++ attributes that convey property names and values to the front-end: @@ -502,11 +520,12 @@ values to the front-end: All of these attributes take a parameter list with the same format. There are always an even number of parameters, where the first half are the property -names and the second half are the property values. The property name is always -a string literal or a `constexpr char *` expression. By convention, property -names that correspond to LLVM IR attributes normally start with the prefix -`"sycl-"` in order to avoid collision with non-SYCL IR attributes, but this is -not a strict requirement. +names and the second half are the property values. (This assumes that the +initial optional parameter is not passed. See below for a description of this +optional parameter.) The property name is always a string literal or a +`constexpr char *` expression. By convention, property names normally start +with the prefix `"sycl-"` in order to avoid collision with non-SYCL IR +attributes, but this is not a strict requirement. The property value can be a literal or `constexpr` expression of the following types: @@ -517,15 +536,29 @@ types: * A boolean type. * A character type. * An enumeration type. +* `nullptr_t` (reserved for the case when a property has no value). All properties require a value when represented in the C++ attribute. If the -SYCL property has no value the header passes the empty string (`""`). +SYCL property has no value the header passes `nullptr`. + +### IR representation as IR attributes + +Properties that are implemented using the following C++ attributes are +represented in LLVM IR as IR attributes: -The front-end converts each value to a string before representing it in LLVM -IR. Integer and floating point values are converted with the same format as -`std::to_string()` would produce. Boolean values are converted to either -`"true"` or `"false"`. Enumeration values are first converted to an integer -and then converted to a string with the same format as `std::to_string()`. +* `[[__sycl_detail__::add_ir_global_variable_attributes()]]` +* `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` +* `[[__sycl_detail__::add_ir_function_attributes()]]` + +When the front-end consumes these C++ attributes and produces IR, each property +name becomes an IR attribute name and each property value becomes the +attribute's value. Because the attribute values must be strings, the front-end +converts each property value to a string. Integer and floating point values +are converted with the same format as `std::to_string()` would produce. +Boolean values are converted to either `"true"` or `"false"`. Enumeration +values are first converted to an integer and then converted to a string with +the same format as `std::to_string()`. The `nullptr` value is converted to an +empty string (`""`). **TODO**: Should we allow property values that are type names? If so, I suppose they would be converted to a string representation of the mangled name? @@ -539,6 +572,94 @@ we do not allow non-fundamental types, how do we represent properties like allow `std::tuple`, where the type of each element is one of the fundamental types listed above. +### IR representation via `@llvm.ptr.annotation` + +Properties that are implemented using +`[[__sycl_detail__::add_ir_member_annotation()]]`, are represented in LLVM IR +as the fifth metadata parameter to the `@llvm.ptr.annotation` intrinsic +function. This parameter is a tuple of metadata values with the following +sequence: + +* Name of the first property +* Value of the first property +* Name of the second property +* Value of the second property +* Etc. + +Since metadata types are not limited to strings, there is no need to convert +the property values to strings. + + +## Filtering properties + +It is sometimes necessary to filter out certain properties so that only a +subset of the properties in a list are represented in IR. There are two +scenarios when this is useful. + +In some cases, a property is used only in the header file itself, and there is +no need to represent that property in LLVM IR. In order to avoid cluttering +the IR with unneeded information, these properties can be "filtered out", so +that the front-end does not generate an IR representation. + +Another case is when a class wants to represent some properties one way in the +IR while representing other properties in another way. For example, a future +version of `accessor` might pass some properties to +`[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` while passing other +properties to `[[__sycl_detail__::add_ir_member_annotation()]]`. Again, the +header wants some way to "filter" the properties, such that some properties are +interpreted as "kernel parameter attributes" while other are interpreted as +"member annotations". + +To handle these cases, each of the following C++ attributes takes an optional +first parameter that is a brace-enclosed list of property names: + +* `[[__sycl_detail__::add_ir_global_variable_attributes()]]` +* `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` +* `[[__sycl_detail__::add_ir_function_attributes()]]` +* `[[__sycl_detail__::add_ir_member_annotation()]]` + +The front-end treats this list as a "pass list", ignoring any property whose +name is not in the list. To illustrate, consider the following example where +`accessor` treats some properties as "kernel parameter attributes" and others +as "member annotations": + +``` +template +class +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_kernel_parameter_attributes( + + // The properties in this list are "kernel parameter attributes". + {"sycl-no-alias", "sycl-foo"}, + + Props::meta_name..., Props::meta_value... + )]] +#endif + accessor> { + T *ptr +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_member_annotation( + + // The properties in this list are "member annotations". + {"sycl-bar"}, + + Props::meta_name..., Props::meta_value... + )]] +#endif + ; + } +``` + ## Representing properties in SPIR-V @@ -638,7 +759,8 @@ of the global variable's decorations. To illustrate: As we noted earlier, a property on a structure member variable is represented in LLVM IR as a call to the intrinsic function `@llvm.ptr.annotation`, where -the annotation string starts with the prefix `"sycl-properties:"`. In order to +the annotation string is `"sycl-properties"` and the properties are represented +as metadata in the fifth parameter to `@llvm.ptr.annotation`. In order to understand how these SYCL properties are translated into SPIR-V, it's useful to review how a normal (i.e. non-SYCL) call to `@llvm.ptr.annotation` is translated. From 4ec288158b945ecf68f098dd82a53e7a263b3014 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Thu, 16 Dec 2021 16:32:42 -0500 Subject: [PATCH 3/7] 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 `` 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 `` is included. --- sycl/doc/CompileTimeProperties.md | 54 ++++++++++++++++++++++--------- 1 file changed, 38 insertions(+), 16 deletions(-) diff --git a/sycl/doc/CompileTimeProperties.md b/sycl/doc/CompileTimeProperties.md index 8b11f4dad9f0b..547801c55bd6b 100644 --- a/sycl/doc/CompileTimeProperties.md +++ b/sycl/doc/CompileTimeProperties.md @@ -89,12 +89,13 @@ class ``` The `[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute has an -even number of parameters. The first half of the parameters are the names of -the properties, and the second half of the parameters are the values for those -properties. Each property has exactly one value, so the property at parameter -position 0 corresponds to the value at position _N / 2_, etc. To illustrate -using the same example as before, the result of the parameter pack expansion -would look like this: +even number of parameters, assuming that the optional "filter list" parameter +is not specified (see below for a description of this parameter). The first +half of the parameters are the names of the properties, and the second half of +the parameters are the values for those properties. Each property has exactly +one value, so the property at parameter position 0 corresponds to the value at +position _N / 2_, etc. To illustrate using the same example as before, the +result of the parameter pack expansion would look like this: ``` namespace sycl::ext::oneapi { @@ -244,12 +245,28 @@ the property value to a string if it is not already a string. [7]: -**TODO**: What happens when a "sycl special class" object is captured as a -kernel argument? The compiler passes each member of the class as a separate -argument. Should the device compiler duplicate the properties on each such -parameter in this case? Or, is it the header's responsibility to add the C++ -attribute to one of the member variables in this case? How does the header -decide which member variable to decorate, though? +**TODO**: There are a number of open issues with this attribute and with the +semantics of properties that are represented as attributes on kernel +arguments. Suppose there are two SYCL types that take properties: _A_ and +_B_. (For example, this could be two specializations of `annotated_ptr`, each +decorated with different properties.) Now suppose the application creates a +struct that contains members with both of these types, and it passes that +struct as a kernel argument. What is the intended semantic? Does the argument +get decorated with the union of the properties on both _A_ and _B_? What if +those properties are mutually exclusive? A similar case exists when the +application creates a struct that inherits from both _A_ and _B_. + +The previous example shows a case when a single kernel argument gets properties +from two (or more) types. However, the opposite can also occur. Certain SYCL +classes are decorated with `__attribute__((sycl_special_class))`, which causes +the compiler to pass each member of that class as a separate kernel argument. +What should happen with the properties that decorate the class? Should the +compiler duplicate the properties on each such kernel argument? Or, maybe it +should be the header file's responsibility not to decorate such a class with +`[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]`, and instead it +should decorate specific member variable(s) with this attribute? How does the +header decide which properties are used to decorate which member variables, +though? ## Properties on kernel functions @@ -344,11 +361,12 @@ string if it is not already a string. [9]: -**TODO**: The intention is to replace the existing member functions like +**NOTE**: The intention is to replace the existing member functions like `handler::kernel_single_task()` with wrapper classes like -`KernelSingleTaskWrapper`. Does this pose any problems? There are comments in -the headers indicating that the front-end recognizes the function -`handler::kernel_single_task()` by name. +`KernelSingleTaskWrapper`. We believe this will not cause problems for the +device compiler front-end because it recognizes kernel functions via the +`__attribute__((sycl_kernel))` attribute, not by the name +`handler::kernel_single_task()`. ## Properties on a non-global variable type @@ -618,6 +636,10 @@ first parameter that is a brace-enclosed list of property names: * `[[__sycl_detail__::add_ir_function_attributes()]]` * `[[__sycl_detail__::add_ir_member_annotation()]]` +Since this brace-enclosed list acts somewhat like an initializer list, the +header must include `` prior to passing this optional first +parameter. + The front-end treats this list as a "pass list", ignoring any property whose name is not in the list. To illustrate, consider the following example where `accessor` treats some properties as "kernel parameter attributes" and others From d0622a60026db7e2ff2d9a69ac8cdfb1ad05e917 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Tue, 11 Jan 2022 16:46:18 -0500 Subject: [PATCH 4/7] Address TODO for kernel arg properties Solve the TODO issues with properties that decorate kernel parameter by: * Move the C++ attribute from the parameter's class to a member variable inside the class. The author of the header file will need to decide with member variable to attach the properties to. * Restrict the C++ attribute, so it is only used to decorate a SYCL "special class". When a value of this type is passed as a kernel parameter, each member variable is passed as a separate parameter to the kernel's function. As a result, there is no ambiguity about which function parameter receives the property. --- sycl/doc/CompileTimeProperties.md | 103 +++++++++++++----------------- 1 file changed, 46 insertions(+), 57 deletions(-) diff --git a/sycl/doc/CompileTimeProperties.md b/sycl/doc/CompileTimeProperties.md index 547801c55bd6b..0e7f072fd43f0 100644 --- a/sycl/doc/CompileTimeProperties.md +++ b/sycl/doc/CompileTimeProperties.md @@ -159,7 +159,7 @@ template > -class accessor {/* ... */}; +class __attribute__((sycl_special_class)) accessor {/* ... */}; } // namespace sycl ``` @@ -176,7 +176,8 @@ accessor acc(buf, cgh, property_list{no_alias_v, foo_v<32>}); As before, the header file represents the properties with an internal C++ attribute, where the initial parameters are the names of the properties and -the subsequent parameters are the property values. +the subsequent parameters are the property values. However, this time the +attribute decorates one of the member variables. ``` namespace sycl { @@ -187,7 +188,7 @@ template > -class accessor {/* ... */}; +class __attribute__((sycl_special_class)) accessor {/* ... */}; // Partial specialization to make PropertyListT visible as a parameter pack // of properties. @@ -197,18 +198,20 @@ template -class +class __attribute__((sycl_special_class)) accessor> { + dataT *ptr #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_kernel_parameter_attributes( Props::meta_name..., Props::meta_value... )]] #endif - accessor> {/*...*/}; + ; +}; } // namespace sycl ``` @@ -218,7 +221,9 @@ Illustrating this with the previous example: ``` namespace sycl { -template class +template +class __attribute__((sycl_special_class)) accessor { + dataT *ptr #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_kernel_parameter_attributes( "sycl-no-alias", // Name of first property @@ -227,47 +232,34 @@ template class 32 // Value of second property )]] #endif - accessor {/* ... */}; + ; +}; } // namespace sycl ``` -As the name of the C++ attribute suggests, the device compiler front-end uses -the attribute only when the decorated type is the type of a kernel argument, +As the name implies, this C++ attribute is only used to decorate a member +variable of a class type that is as SYCL "special class" (i.e. a class that is +decorated with `__attribute__((sycl_special_class))`). The device compiler +front-end ignores the attribute when it is used in any other syntactic +position. + +The device compiler front-end uses this attribute only when the class type +containing the decorated member variable is the type of a kernel argument, and it silently ignores the attribute when the class is used in any other way. -When the device compiler front-end creates a kernel argument in this way, it -adds one LLVM IR attribute to the kernel function's parameter for each property -in the list. For example, this can be done by calling +When the front-end creates a kernel argument from a SYCL "special class", it +passes each member variable of the class as a separate kernel argument. If the +member variable is decorated with +`[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]`, the front-end adds +one LLVM IR attribute to the kernel function's parameter for each property in +the list. For example, this can be done by calling [`Function::addParamAttrs(unsigned ArgNo, const AttrBuilder &)`][7]. As before, the IR attributes are added as strings, so the front-end must convert the property value to a string if it is not already a string. [7]: -**TODO**: There are a number of open issues with this attribute and with the -semantics of properties that are represented as attributes on kernel -arguments. Suppose there are two SYCL types that take properties: _A_ and -_B_. (For example, this could be two specializations of `annotated_ptr`, each -decorated with different properties.) Now suppose the application creates a -struct that contains members with both of these types, and it passes that -struct as a kernel argument. What is the intended semantic? Does the argument -get decorated with the union of the properties on both _A_ and _B_? What if -those properties are mutually exclusive? A similar case exists when the -application creates a struct that inherits from both _A_ and _B_. - -The previous example shows a case when a single kernel argument gets properties -from two (or more) types. However, the opposite can also occur. Certain SYCL -classes are decorated with `__attribute__((sycl_special_class))`, which causes -the compiler to pass each member of that class as a separate kernel argument. -What should happen with the properties that decorate the class? Should the -compiler duplicate the properties on each such kernel argument? Or, maybe it -should be the header file's responsibility not to decorate such a class with -`[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]`, and instead it -should decorate specific member variable(s) with this attribute? How does the -header decide which properties are used to decorate which member variables, -though? - ## Properties on kernel functions @@ -402,8 +394,8 @@ void foo(int *p) { } ``` -We again implement the property list in the header via a C++ attribute, though -this time the attribute decorates a member variable of the class: +We again implement the property list in the header via a C++ attribute, where +the attribute decorates a member variable of the class: ``` namespace sycl::ext::oneapi { @@ -652,24 +644,21 @@ template -class +class __attribute__((sycl_special_class)) accessor> { + T *ptr #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_kernel_parameter_attributes( + [[__sycl_detail__::add_ir_kernel_parameter_attributes( - // The properties in this list are "kernel parameter attributes". - {"sycl-no-alias", "sycl-foo"}, + // The properties in this list are "kernel parameter attributes". + {"sycl-no-alias", "sycl-foo"}, - Props::meta_name..., Props::meta_value... - )]] -#endif - accessor> { - T *ptr -#ifdef __SYCL_DEVICE_ONLY__ + Props::meta_name..., Props::meta_value... + )]] [[__sycl_detail__::add_ir_member_annotation( // The properties in this list are "member annotations". From cf066942c65040b6cb3c975b9be1318318a253f4 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Wed, 12 Jan 2022 11:40:51 -0500 Subject: [PATCH 5/7] Add property design doc to index --- sycl/doc/index.rst | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/index.rst b/sycl/doc/index.rst index 8945b41423953..844e0c1c3be7b 100644 --- a/sycl/doc/index.rst +++ b/sycl/doc/index.rst @@ -41,6 +41,7 @@ Developing oneAPI DPC++ Compiler SYCLInstrumentationUsingXPTI ITTAnnotations DeviceGlobal + CompileTimeProperties Development BKMs ~~~~~~~~~~~~~~~~ From 403e6c964c4539197fc186aee9a9ce2ff271ebe7 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Thu, 13 Jan 2022 10:36:43 -0500 Subject: [PATCH 6/7] Clarify wording about C++ attributes Clarify the description about the changes to the header file to state exactly which C++ attribute is added. --- sycl/doc/CompileTimeProperties.md | 24 +++++++++++++++--------- 1 file changed, 15 insertions(+), 9 deletions(-) diff --git a/sycl/doc/CompileTimeProperties.md b/sycl/doc/CompileTimeProperties.md index 0e7f072fd43f0..4c8c66585481a 100644 --- a/sycl/doc/CompileTimeProperties.md +++ b/sycl/doc/CompileTimeProperties.md @@ -174,10 +174,11 @@ using sycl::ext::oneapi; accessor acc(buf, cgh, property_list{no_alias_v, foo_v<32>}); ``` -As before, the header file represents the properties with an internal C++ -attribute, where the initial parameters are the names of the properties and -the subsequent parameters are the property values. However, this time the -attribute decorates one of the member variables. +The implementation in the header file is similar to the previous case. The +C++ attribute `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` +decorates one of the member variables of the class, and the parameters to this +attribute represent the properties. As before, the initial parameters are the +names of the properties and the subsequent parameters are the property values. ``` namespace sycl { @@ -315,9 +316,12 @@ void foo(handler &cgh) { } ``` -Internally, the headers lower both cases to a wrapper class that expresses the -properties as an internal C++ attribute, and the `operator()` of this class -becomes the "top level" kernel function that is recognized by the front-end. +Internally, the header lowers both cases to a wrapper class which defines +`operator()`, and that operator function becomes the "top level" kernel +function that is recognized by the front-end. The definition of this operator +is decorated with the C++ attribute +`[[__sycl_detail__::add_ir_function_attributes()]]`, and the parameters to this +attribute represent the properties. ``` template @@ -394,8 +398,10 @@ void foo(int *p) { } ``` -We again implement the property list in the header via a C++ attribute, where -the attribute decorates a member variable of the class: +We again use a C++ attribute to represent the properties in the header. The +attribute `[[__sycl_detail__::add_ir_member_annotation()]]` decorates one of +the member variables of the class, and the parameters to this attribute +represent the properties. ``` namespace sycl::ext::oneapi { From 9ad46514427b6e99783f955e6092cccaafacdac2 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Thu, 13 Jan 2022 13:29:26 -0500 Subject: [PATCH 7/7] Rename C++ attributes --- sycl/doc/CompileTimeProperties.md | 64 +++++++++++++++---------------- 1 file changed, 32 insertions(+), 32 deletions(-) diff --git a/sycl/doc/CompileTimeProperties.md b/sycl/doc/CompileTimeProperties.md index 4c8c66585481a..60b89825e7579 100644 --- a/sycl/doc/CompileTimeProperties.md +++ b/sycl/doc/CompileTimeProperties.md @@ -65,7 +65,7 @@ device_global class #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_global_variable_attributes( + [[__sycl_detail__::add_ir_attributes_global_variable( Props::meta_name..., Props::meta_value... )]] #endif @@ -88,7 +88,7 @@ class } // namespace sycl::ext::oneapi ``` -The `[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute has an +The `[[__sycl_detail__::add_ir_attributes_global_variable()]]` attribute has an even number of parameters, assuming that the optional "filter list" parameter is not specified (see below for a description of this parameter). The first half of the parameters are the names of the properties, and the second half of @@ -102,7 +102,7 @@ namespace sycl::ext::oneapi { template class #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_global_variable_attributes( + [[__sycl_detail__::add_ir_attributes_global_variable( "sycl-device-image-scope", // Name of first property "sycl-host-access", // Name of second property nullptr, // First property has no parameter @@ -115,7 +115,7 @@ template class ``` The device compiler only uses the -`[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute when the +`[[__sycl_detail__::add_ir_attributes_global_variable()]]` attribute when the decorated type is used to create an [LLVM IR global variable][3] and the global variable's type is either: @@ -175,7 +175,7 @@ accessor acc(buf, cgh, property_list{no_alias_v, foo_v<32>}); ``` The implementation in the header file is similar to the previous case. The -C++ attribute `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` +C++ attribute `[[__sycl_detail__::add_ir_attributes_kernel_parameter()]]` decorates one of the member variables of the class, and the parameters to this attribute represent the properties. As before, the initial parameters are the names of the properties and the subsequent parameters are the property values. @@ -207,7 +207,7 @@ class __attribute__((sycl_special_class)) accessor> { dataT *ptr #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_kernel_parameter_attributes( + [[__sycl_detail__::add_ir_attributes_kernel_parameter( Props::meta_name..., Props::meta_value... )]] #endif @@ -226,7 +226,7 @@ template class __attribute__((sycl_special_class)) accessor { dataT *ptr #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_kernel_parameter_attributes( + [[__sycl_detail__::add_ir_attributes_kernel_parameter( "sycl-no-alias", // Name of first property "sycl-foo", // Name of second property nullptr, // First property has no parameter @@ -252,7 +252,7 @@ and it silently ignores the attribute when the class is used in any other way. When the front-end creates a kernel argument from a SYCL "special class", it passes each member variable of the class as a separate kernel argument. If the member variable is decorated with -`[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]`, the front-end adds +`[[__sycl_detail__::add_ir_attributes_kernel_parameter()]]`, the front-end adds one LLVM IR attribute to the kernel function's parameter for each property in the list. For example, this can be done by calling [`Function::addParamAttrs(unsigned ArgNo, const AttrBuilder &)`][7]. As @@ -320,7 +320,7 @@ Internally, the header lowers both cases to a wrapper class which defines `operator()`, and that operator function becomes the "top level" kernel function that is recognized by the front-end. The definition of this operator is decorated with the C++ attribute -`[[__sycl_detail__::add_ir_function_attributes()]]`, and the parameters to this +`[[__sycl_detail__::add_ir_attributes_function()]]`, and the parameters to this attribute represent the properties. ``` @@ -338,7 +338,7 @@ class KernelSingleTaskWrapper> { #ifdef __SYCL_DEVICE_ONLY__ __attribute__((sycl_kernel)) - [[__sycl_detail__::add_ir_function_attributes( + [[__sycl_detail__::add_ir_attributes_function( Props::meta_name..., Props::meta_value... )]] #endif @@ -347,7 +347,7 @@ class KernelSingleTaskWrapper> { ``` Although the DPC++ headers only use the -`[[__sycl_detail__::add_ir_function_attributes()]]` attribute on the definition +`[[__sycl_detail__::add_ir_attributes_function()]]` attribute on the definition of a kernel function as shown above, the front-end recognizes it for any function definition. The front-end adds one LLVM IR function attribute for each property in the list. For example, this can be done by calling @@ -399,7 +399,7 @@ void foo(int *p) { ``` We again use a C++ attribute to represent the properties in the header. The -attribute `[[__sycl_detail__::add_ir_member_annotation()]]` decorates one of +attribute `[[__sycl_detail__::add_ir_annotations_member()]]` decorates one of the member variables of the class, and the parameters to this attribute represent the properties. @@ -415,7 +415,7 @@ template class annotated_ptr> { T *ptr #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_member_annotation( + [[__sycl_detail__::add_ir_annotations_member( Props::meta_name..., Props::meta_value... )]] #endif @@ -441,7 +441,7 @@ template class annotated_ptr> { T *ptr #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_member_annotation( + [[__sycl_detail__::add_ir_annotations_member( "sycl-foo", // Name of first property "sycl-bar", // Name of second property nullptr, // First property has no parameter @@ -501,7 +501,7 @@ define void @foo(i32* %ptr) { ``` The front-end encodes the properties from the C++ attribute -`[[__sycl_detail__::add_ir_member_annotation()]]` into the +`[[__sycl_detail__::add_ir_annotations_member()]]` into the `@llvm.ptr.annotation` call as follows: * The first parameter to `@llvm.ptr.annotation` is the pointer to annotate (as @@ -529,10 +529,10 @@ to perform these optimizations. As noted above, there are several C++ attributes that convey property names and values to the front-end: -* `[[__sycl_detail__::add_ir_global_variable_attributes()]]` -* `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` -* `[[__sycl_detail__::add_ir_function_attributes()]]` -* `[[__sycl_detail__::add_ir_member_annotation()]]` +* `[[__sycl_detail__::add_ir_attributes_global_variable()]]` +* `[[__sycl_detail__::add_ir_attributes_kernel_parameter()]]` +* `[[__sycl_detail__::add_ir_attributes_function()]]` +* `[[__sycl_detail__::add_ir_annotations_member()]]` All of these attributes take a parameter list with the same format. There are always an even number of parameters, where the first half are the property @@ -562,9 +562,9 @@ SYCL property has no value the header passes `nullptr`. Properties that are implemented using the following C++ attributes are represented in LLVM IR as IR attributes: -* `[[__sycl_detail__::add_ir_global_variable_attributes()]]` -* `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` -* `[[__sycl_detail__::add_ir_function_attributes()]]` +* `[[__sycl_detail__::add_ir_attributes_global_variable()]]` +* `[[__sycl_detail__::add_ir_attributes_kernel_parameter()]]` +* `[[__sycl_detail__::add_ir_attributes_function()]]` When the front-end consumes these C++ attributes and produces IR, each property name becomes an IR attribute name and each property value becomes the @@ -591,7 +591,7 @@ types listed above. ### IR representation via `@llvm.ptr.annotation` Properties that are implemented using -`[[__sycl_detail__::add_ir_member_annotation()]]`, are represented in LLVM IR +`[[__sycl_detail__::add_ir_annotations_member()]]`, are represented in LLVM IR as the fifth metadata parameter to the `@llvm.ptr.annotation` intrinsic function. This parameter is a tuple of metadata values with the following sequence: @@ -620,8 +620,8 @@ that the front-end does not generate an IR representation. Another case is when a class wants to represent some properties one way in the IR while representing other properties in another way. For example, a future version of `accessor` might pass some properties to -`[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` while passing other -properties to `[[__sycl_detail__::add_ir_member_annotation()]]`. Again, the +`[[__sycl_detail__::add_ir_attributes_kernel_parameter()]]` while passing other +properties to `[[__sycl_detail__::add_ir_annotations_member()]]`. Again, the header wants some way to "filter" the properties, such that some properties are interpreted as "kernel parameter attributes" while other are interpreted as "member annotations". @@ -629,10 +629,10 @@ interpreted as "kernel parameter attributes" while other are interpreted as To handle these cases, each of the following C++ attributes takes an optional first parameter that is a brace-enclosed list of property names: -* `[[__sycl_detail__::add_ir_global_variable_attributes()]]` -* `[[__sycl_detail__::add_ir_kernel_parameter_attributes()]]` -* `[[__sycl_detail__::add_ir_function_attributes()]]` -* `[[__sycl_detail__::add_ir_member_annotation()]]` +* `[[__sycl_detail__::add_ir_attributes_global_variable()]]` +* `[[__sycl_detail__::add_ir_attributes_kernel_parameter()]]` +* `[[__sycl_detail__::add_ir_attributes_function()]]` +* `[[__sycl_detail__::add_ir_annotations_member()]]` Since this brace-enclosed list acts somewhat like an initializer list, the header must include `` prior to passing this optional first @@ -658,14 +658,14 @@ class __attribute__((sycl_special_class)) accessor> { T *ptr #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_kernel_parameter_attributes( + [[__sycl_detail__::add_ir_attributes_kernel_parameter( // The properties in this list are "kernel parameter attributes". {"sycl-no-alias", "sycl-foo"}, Props::meta_name..., Props::meta_value... )]] - [[__sycl_detail__::add_ir_member_annotation( + [[__sycl_detail__::add_ir_annotations_member( // The properties in this list are "member annotations". {"sycl-bar"},