SYCL binary property sets

To communicate information about SYCL binaries to the SYCL runtime, the implementation produces sets of properties. The intention of this design document is to describe the structure of the property sets and define the representation and meaning of pre-defined property set names.

Property sets structure

A property set consists of a reserved name, enclosed in square brackets, followed by a series of string key and value pairs. The set name and each entry in the set are separated by a newline.

The string key and value pairs have the following format:

<string key>=<value type>|<value>

The value type is a string and the value of it has the following meaning for the corresponding value:

Value type	Description
“1”	The value is a 32 bit integer.
“2”	The value is a base64 encoded byte array.

Note: Whitespaces are not ignored and are treated like any other characters. As such, some_key=1|1 is not the same as some_key = 1 | 1 and [some property set] is not the same as [ some property set ].

Property sets

This section describes the known property sets.

[SYCL/specialization constants]

Key: Specialization constant name.

Value type: Byte array. (“2”)

Value: Information about the specialization constant with the following fields:

// Encodes ID of a scalar specialization constants which is a leaf of some
// composite specialization constant.
unsigned ID;
// Encodes offset from the beginning of composite, where scalar resides, i.e.
// location of the scalar value within a byte-array containing the whole
// composite specialization constant. If descriptor is used to represent a
// whole scalar specialization constant instead of an element of a composite,
// this field should be contain zero.
unsigned Offset;
// Encodes size of scalar specialization constant.
unsigned Size;

See also SYCL2020-SpecializationConstants.md.

[SYCL/specialization constants default values]

Key: Specialization constant name.

Value type: Byte array. (“2”)

Value: Byte representation of the default value for the specialization constant.

See also SYCL2020-SpecializationConstants.md.

[SYCL/devicelib req mask]

Key: At most one entry with “DeviceLibReqMask”.

Value type: 32 bit integer. (“1”)

Value: A bitmask of which device libraries the binary uses.

Notes:

1. If this property set is missing, no device libraries are used by the binary.

[SYCL/kernel param opt]

Key: Kernel name.

Value type: Byte array. (“2”)

Value: A bitmask identifying the arguments of the kernel that have been removed by the dead-argument-elimination optimization pass.

Notes:

1. If no entry is present for a given kernel in the binary, no arguments have been eliminated.

2. If this property set is missing, no kernels in the binary have any eliminated arguments.

[SYCL/program metadata]

Program metadata properties:

Key	Value type	Value
kernel + “@reqd_work_group_size”	Byte array. (“2”)	Specifies the required work-group size for the kernel identified by the name kernel.
kernel + “@work_group_num_dim”	Byte array. (“2”)	Specifies the work-group dimensionality of the kernel identified by the name kernel.
kernel + “@max_work_group_size”	Byte array. (“2”)	Specifies the max work-group size for the kernel identified by the name kernel.
kernel + “@max_linear_work_group_size”	Byte array. (“2”)	Specifies the max linear work-group size for the kernel identified by the name kernel.
variable + “@global_id_mapping”	Byte array. (“2”)	Specifies the mapping between the global variable with unique identifier variable and its name in the binary.

[SYCL/misc properties]

Miscellaneous properties:

Key	Value type	Value
“isEsimdImage”	32 bit integer. (“1”)	1 if the image is ESIMD and 0 or missing otherwise.
“sycl-register-alloc-mode”	32 bit integer. (“1”)	The register allocation mode: 2 for large and 0 or missing for automatic.
“sycl-grf-size”	32 bit integer. (“1”)	The GRF size. Automatic if 0 or missing.
“optLevel”	32 bit integer. (“1”)	Optimization level, corresponding to the -O option used during compilation.
“sanUsed”	Byte array. (“2”)	Specifying if address sanitization (“asan”) or memory sanitization (“msan”) is used. Missing if neither is used.
“specConstsReplacedWithDefault”	32 bit integer. (“1”)	1 if the specialization constants have been replaced by their default values and 0 or missing otherwise.

[SYCL/assert used]

Key: Kernel name.

Value type: 32 bit integer. (“1”)

Value: 1 if the kernel uses assertions and 0 or missing otherwise.

[SYCL/exported symbols]

Key: Symbol name.

Value type: 32 bit integer. (“1”)

Value: 1 if the symbol is exported by the binary and 0 or missing otherwise.

See also SharedLibraries.md.

[SYCL/imported symbols]

Key: Symbol name.

Value type: 32 bit integer. (“1”)

Value: 1 if the symbol is imported by the binary and 0 or missing otherwise.

See also SharedLibraries.md.

[SYCL/device globals]

Key: Device global variable name.

Value type: Byte array. (“2”)

Value: Information about the device global variable with the following fields:

  // Encodes size of the underlying type T of the device global variable.
  uint32_t Size;

  // Either 1 (true) or 0 (false), telling whether the device global variable
  // was declared with the device_image_scope property.
  // We use uint32_t for a boolean value to eliminate padding after the field
  // and suppress false positive reports from MemorySanitizer.
  uint32_t DeviceImageScope;

Notes:

1. If this property set is missing, the binary does not contain any device global variables.

See also DeviceGlobal.md.

[SYCL/device requirements]

Set of device requirements for the entire module:

Key	Value type	Value
“aspects”	Byte array. (“2”)	A collection of 32 bit integers representing the SYCL aspects used. These correspond 1:1 with the enum values of sycl::aspect.
“fixed_target”	Byte array. (“2”)	The string literals specified in -fsycl-fixed-targets.
“reqd_work_group_size_uint64_t”	Byte array. (“2”)	At most three 64 bit unsigned integers representing the required work-group size. If this entry is missing, there is no work-group size requirement.
“joint_matrix”	Byte array. (“2”)	A string containing a semi-colon-separated list of comma-separated descriptors for used matrices. The descriptors in the order they appear are: sycl-joint-matrix-type sycl-joint-matrix-use sycl-joint-matrix-rows sycl-joint-matrix-cols
“joint_matrix_mad”	Byte array. (“2”)	A string containing a semi-colon-separated list of comma-separated descriptors for used matrix MAD operations. The descriptors in the order they appear are: sycl-joint-matrix-mad-type-A sycl-joint-matrix-mad-type-B sycl-joint-matrix-mad-type-C sycl-joint-matrix-mad-type-D sycl-joint-matrix-mad-size-M sycl-joint-matrix-mad-size-K sycl-joint-matrix-mad-size-N
“reqd_sub_group_size”	Byte array. (“2”)	At most three 32 bit unsigned integers representing the required sub-group size. If this entry is missing, there is no sub-group size requirement.
“work_group_num_dim”	Byte array. (“2”)	At most three 32 bit unsigned integers representing the work-group dimensionality. If this entry is missing, there is no specified work-group dimensionality.

See also OptionalDeviceFeatures.md.

[SYCL/host pipes]

Key: Host pipe variable name.

Value type: Byte array. (“2”)

Value: Information about the host pipe variable with the following fields:

  // Encodes size of the underlying type T of the host pipe variable.
  uint32_t Size;

Notes:

1. If this property set is missing, the binary does not contain any host pipe variables.

[SYCL/virtual functions]

Set of information about virtual function usage in the module.

Key	Value type	Value
“virtual-functions-set”	Byte array. (“2”)	A string identifying the set of virtual functions contained in the module. If this is missing, the module does not contain any virtual function sets.
“uses-virtual-functions-set”	Byte array. (“2”)	A string containing a comma-separated list of sets of virtual functions used by the module. If this is missing, the module does not use any virtual function sets.

[SYCL/implicit local arg]

Key: Kernel name.

Value type: 32 bit integer. (“1”)

Value: Index of the implicit local memory argument.

Notes:

1. If no entry is present for a given kernel in the binary, the kernel does not have an implicit local memory argument.

2. If this property set is missing, no kernels in the binary have an implicit local memory argument.

[SYCL/registered kernels]

Key: “Registered” kernel name.

Value type: Byte array. (“2”)

Value: The name of the kernel corresponding to the registered kernel name.

Notes:

1. If this property set is missing, the binary does not have any registered kernel names.

[SYCLBIN/global metadata]

Set of global information about a SYCLBIN file.

Key	Value type	Value
“state”	32 bit integer. (“1”)	Integer representation of one of the possible states of the file, corresponding to the sycl::bundle_state enum. It must be one of the following: sycl::bundle_state::input sycl::bundle_state::object sycl::bundle_state::executable

[SYCLBIN/ir module metadata]

Set of information about an IR module in a SYCLBIN file.

Key	Value type	Value
“type”	32 bit integer. (“1”)	Integer representation of one of the pre-defined IR types. It must be one of the following: SPIR-V PTX AMDGCN
“target”	Byte array. (“2”)	A string representing the architecture of the binary, corresponding to the value of -fsycl-targets option used when compiling this binary. This may be missing if no part of -fsycl-targets was used during the compilation of this binary or if -fsycl-targets was not used at all.

[SYCLBIN/native device code image metadata]

Set of information about an native device code image in a SYCLBIN file.

Key	Value type	Value
“arch”	Byte array. (“2”)	A string representing the architecture of the binary, corresponding to the value of -fsycl-targets option used when compiling this binary.