Implementation design for compile time constant properties¶
This document describes the implementation design for the DPC++ extension sycl_ext_oneapi_properties, 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.
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 extension:
namespace sycl::ext::oneapi {
template <typename T, typename PropertyListT = experimental::empty_properties_t>
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<int, decltype(properties{device_image_scope, host_access_read})>
dm1;
The header file represents these properties with an internal C++ attribute
named [[__sycl_detail__::add_ir_attributes_global_variable()]] whose value
is a list that is created through a template parameter pack expansion:
namespace sycl::ext::oneapi {
template <typename T, typename PropertyListT = experimental::empty_properties_t>
class device_global {/*...*/};
// Partial specialization to make PropertyListT visible as a parameter pack
// of properties.
template <typename T, typename ...Props>
class
#ifdef __SYCL_DEVICE_ONLY__
[[__sycl_detail__::add_ir_attributes_global_variable(
Props::meta_name..., Props::meta_value...
)]]
#endif
device_global<T, properties<Props...>> {/*...*/};
} // namespace sycl::ext::oneapi
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
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 {
template </* ... */> class
#ifdef __SYCL_DEVICE_ONLY__
[[__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
"read" // Value of second property
)]]
#endif
device_global</* ... */> {/*...*/};
} // namespace sycl::ext::oneapi
The device compiler only uses the
[[__sycl_detail__::add_ir_attributes_global_variable()]] attribute when the
decorated type is used to create an LLVM IR global variable 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.
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). If the property
value is not already a string, it converts it to a string as described in
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.
Properties on kernel arguments¶
Another use of compile-time properties is with types that are used to define kernel arguments. For example, the sycl_ext_oneapi_accessor_properties extension could be redesigned to use compile-time properties. Such a redesign might look like:
namespace sycl {
template <typename dataT,
int dimensions,
access::mode accessmode,
access::target accessTarget,
access::placeholder isPlaceholder,
typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
class __attribute__((sycl_special_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, properties{no_alias, foo<32>});
In the headers the C++ attribute
[[__sycl_detail__::add_ir_attributes_kernel_parameter()]] is used to decorate
parameters of the __init member function in the corresponding
sycl_special_class decorated class. As before, the initial parameters are the
names of the properties and the subsequent parameters are the property values.
namespace sycl {
template <typename dataT,
int dimensions,
access::mode accessmode,
access::target accessTarget,
access::placeholder isPlaceholder,
typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
class __attribute__((sycl_special_class)) accessor {/* ... */};
// Partial specialization to make PropertyListT visible as a parameter pack
// of properties.
template <typename dataT,
int dimensions,
access::mode accessmode,
access::target accessTarget,
access::placeholder isPlaceholder,
typename ...Props>
class __attribute__((sycl_special_class)) accessor<dataT,
dimensions,
accessmode,
accessTarget,
isPlaceholder,
properties<Props...>> {
dataT *ptr;
#ifdef __SYCL_DEVICE_ONLY__
void __init(
[[__sycl_detail__::add_ir_attributes_kernel_parameter(
Props::meta_name..., Props::meta_value...
)]]
dataT *_ptr) {
ptr = _ptr;
}
#endif
};
} // namespace sycl
Illustrating this with the previous example:
namespace sycl {
template </* ... */>
class __attribute__((sycl_special_class)) accessor</* ... */> {
dataT *ptr;
#ifdef __SYCL_DEVICE_ONLY__
void __init(
[[__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
32 // Value of second property
)]]
dataT *_ptr) {
ptr = _ptr;
}
#endif
};
} // namespace sycl
As the name implies, this C++ attribute is only used to decorate parameters of
the __init member function 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.
When the front-end creates a kernel argument from a SYCL “special class”, it
copies all parameters of the __init member function to the corresponding
kernel function. If a copied parameter is decorated with
[[__sycl_detail__::add_ir_attributes_kernel_parameter()]], the front-end adds
one LLVM IR attribute to the resulting kernel function parameter for each
property in the list. For example, this can be done by calling
Function::addParamAttrs(unsigned ArgNo, const AttrBuilder &). 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.
Properties on kernel functions¶
Compile-time properties can also be used to decorate kernel functions as
proposed in the sycl_ext_oneapi_kernel_properties extension. There
are two ways the application can specify these properties. The first is by
passing a properties parameter to the function that submits the kernel:
namespace sycl {
class handler {
template <typename KernelName, typename KernelType, typename PropertyListT>
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(
properties{sub_group_size<32>, device_has<aspect::fp16>},
[=] {/* ... */});
}
The second way an application can specify kernel properties is by adding a
member function named get(sycl::ext::oneapi::properties_tag) to a named
kernel function object:
using sycl;
using sycl::ext::oneapi;
class MyKernel {
public:
void operator()() {/* ... */}
auto get(properties_tag) {
return properties{sub_group_size<32>, device_has<aspect::fp16>};
}
};
void foo(handler &cgh) {
MyKernel k;
cgh.single_task(k);
}
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_attributes_function()]], and the parameters to this
attribute represent the properties.
template<typename KernelType, typename PropertyListT>
class KernelSingleTaskWrapper;
// Partial specialization to make PropertyListT visible as a parameter pack
// of properties.
template<typename KernelType, typename ...Props>
class KernelSingleTaskWrapper<KernelType, properties<Props...>> {
KernelType k;
public:
KernelSingleTaskWrapper(KernelType k) : k(k) {}
#ifdef __SYCL_DEVICE_ONLY__
[[clang::sycl_kernel]]
[[__sycl_detail__::add_ir_attributes_function(
Props::meta_name..., Props::meta_value...
)]]
#endif
void operator()() const {k();}
};
Although the DPC++ headers only use the
[[__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
Function::addFnAttr(StringRef, StringRef). 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.
NOTE: The intention is to replace the existing member functions like
handler::kernel_single_task() with wrapper classes like
KernelSingleTaskWrapper. We believe this will not cause problems for the
device compiler front-end because it recognizes kernel functions via the
[[clang::sycl_kernel]] attribute, not by the name
handler::kernel_single_task().
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 <typename T, typename PropertyListT = experimental::empty_properties_t>
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<int, decltype(properties{foo, bar<32>})> aptr(p);
}
We again use a C++ attribute to represent the properties in the header. The
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.
namespace sycl::ext::oneapi {
template <typename T, typename PropertyListT = experimental::empty_properties_t>
class annotated_ptr;
// Partial specialization to make PropertyListT visible as a parameter pack
// of properties.
template <typename T, typename ...Props>
class annotated_ptr<T, properties<Props...>> {
T *ptr
#ifdef __SYCL_DEVICE_ONLY__
[[__sycl_detail__::add_ir_annotations_member(
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 <typename T, typename PropertyListT = experimental::empty_properties_t>
class annotated_ptr;
// Partial specialization to make PropertyListT visible as a parameter pack
// of properties.
template <typename T, typename ...Props>
class annotated_ptr<T, properties<Props...>> {
T *ptr
#ifdef __SYCL_DEVICE_ONLY__
[[__sycl_detail__::add_ir_annotations_member(
"sycl-foo", // Name of first property
"sycl-bar", // Name of second property
nullptr, // 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 that annotates the pointer to that member
variables, similar to the way the existing [[clang::annotate()]] attribute
works.
The front-end encodes the properties from the C++ attribute
[[__sycl_detail__::add_ir_annotations_member()]] into the
@llvm.ptr.annotation call as follows:
The first parameter to
@llvm.ptr.annotationis 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 pointer to a constant global variable. The type of this variable is an anonymous structure. The first field of the structure is a pointer to a string literal representing the name of the first property. The second field of the structure is a pointer to a string literal representing the value of the first property. The third field of the structure is a pointer to a string literal representing the name of the second property, etc. Since each property has exactly one value, this tuple has an even number of elements. Pointers to property value strings may be a null-pointer, signalling a property without a value.
The resulting LLVM IR for the previous example would be:
@.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"
@.str.2 = private unnamed_addr constant [9 x i8] c"sycl-foo\00",
section "llvm.metadata"
@.str.3 = private unnamed_addr constant [9 x i8] c"sycl-bar\00",
section "llvm.metadata"
@.str.4 = private unnamed_addr constant [3 x i8] c"32\00",
section "llvm.metadata"
@.args = private unnamed_addr constant { [9 x i8]*, i8*, [9 x i8]*, [3 x i8]* }
{
[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"
[3 x i8]* @.str.4 ; 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* 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]*, [3 x i8]* }* @.args to i8*))
%3 = bitcast i8* %2 to i32**
store i32* %ptr, i32** %3
ret void
}
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 and in IR¶
As noted above, there are several C++ attributes that convey property names and values to the front-end:
[[__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
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:
const char *.An integer type.
A floating point type.
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 nullptr.
IR representation as IR attributes¶
Properties that are implemented using the following C++ attributes are represented in LLVM IR as IR 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
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?
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.
IR representation via @llvm.ptr.annotation¶
Properties that are implemented using
[[__sycl_detail__::add_ir_annotations_member()]], are represented in LLVM IR
as the fifth parameter to the @llvm.ptr.annotation intrinsic function. This
parameter is a pointer to a global variable with fields corresponding to the
names and values of the properties in the following sequence:
Name of the first property
Value of the first property
Name of the second property
Value of the second property
Etc.
Every field in the global variable pointed to by this parameter are string
literals in seperate global variables. Property values are converted to strings
in the same way as described above, except that the nullptr value and the
empty string ("") is represented as null in the global variable field.
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_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”.
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_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 <initializer_list> 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
as “member annotations”:
template <typename dataT,
int dimensions,
access::mode accessmode,
access::target accessTarget,
access::placeholder isPlaceholder,
typename ...Props>
class __attribute__((sycl_special_class)) accessor<dataT,
dimensions,
accessmode,
accessTarget,
isPlaceholder,
properties<Props...>> {
T *ptr
#ifdef __SYCL_DEVICE_ONLY__
[[__sycl_detail__::add_ir_annotations_member(
// The properties in this list are "member annotations".
{"sycl-bar"},
Props::meta_name..., Props::meta_value...
)]]
#endif
;
#ifdef __SYCL_DEVICE_ONLY__
void __init(
[[__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...
)]]
dataT *_ptr) {
ptr = _ptr;
}
#endif
}
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 to translate from LLVM IR to SPIR-V,
and that tool defines idiomatic LLVM IR 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.
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 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.
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.
An exception to this is for a selection of FPGA-related decorations. If these are supported during translation from LLVM IR to SPIR-V the corresponding decorations will be generated, and otherwise it will fall back to creating a single UserSemantic decoration. In general these decorations occur in the annotation string as a series of {X} and {X:Y} where X is a reserved name and Y is one or more words and numbers separated by a comma (,) or a colon (:), depending on the decoration.
As such we propose an extension to this functionality with the following changes:
To bring it in line with the format of the metadata decorations, the parsing of these decorations should allow the use of SPIR-V decoration identifiers rather than reserved names. With this there need not be any agreement between the translator and LLVM IR producer, as the identifiers are specified by the SPIR-V specification.
For decorations parsed with decoration identifiers, only the comma delimiter is valid for separating decoration values.
In addition to words and numbers, string literals enclosed by quotation marks are allowed as decoration values. No escapes are planned for this, so all symbols between starting quotation mark and ending quotation mark are considered part of the string literal.
When a member variable property needs to be represented in SPIR-V, the
sycl-post-link tool converts the @llvm.ptr.annotation intrinsic call
produced by [[__sycl_detail__::add_ir_annotations_member()]] into another
@llvm.ptr.annotation intrinsic call using this format. For example:
; Contains decorations:
; * 7744 with no value.
; * 7745 with 20 and "str 1" as the values.
@.str = private unnamed_addr constant [24 x i8] c"{7744}{7745:20,\22str 1\22}\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) {
...
; %0 points to the annotated member field.
%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* null)
...
}
NOTE: To allow backwards compatibility with the old format, reverse translation of decorations will produce a decorations in the annotation string following the old format if the decoration had a reserved name.