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 = properties<>>
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 = properties<>>
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::properties<>>
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::properties<>>
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 = properties<>>
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 = properties<>>
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 = properties<>>
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.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 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.