[SYCL][Docs] Add SYCLBIN feature and format design document (#16872)

This commit adds a design document detailing the SYCLBIN binary format
for representing SYCL device kernel binaries to be loaded dynamically at
runtime. Additionally, the design document details how this is to be
handled by the SYCL runtime, driver and clang tooling.

As the design of SYCLBIN files relies heavily on the property sets, this
PR also adds documentation to the existing property set functionality.

---------

Signed-off-by: Larsen, Steffen <[email protected]>
Co-authored-by: Michael Toguchi <[email protected]>
Co-authored-by: Greg Lueck <[email protected]>
Co-authored-by: Alexey Sachkov <[email protected]>

Author: 4 people

Diff for: sycl/doc/design/PropertySets.md

@@ -0,0 +1,296 @@
+# 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: 
+
+```c++
+// 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](./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](./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](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](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: 
+
+```c++
+  // 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](./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: <ul><li>sycl-joint-matrix-type</li><li>sycl-joint-matrix-use</li><li>sycl-joint-matrix-rows</li><li>sycl-joint-matrix-cols</li></ul> |
+| "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: <ul><li>sycl-joint-matrix-mad-type-A</li><li>sycl-joint-matrix-mad-type-B</li><li>sycl-joint-matrix-mad-type-C</li><li>sycl-joint-matrix-mad-type-D</li><li>sycl-joint-matrix-mad-size-M</li><li>sycl-joint-matrix-mad-size-K</li><li>sycl-joint-matrix-mad-size-N</li></ul> |
+| "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](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: 
+
+```c++
+  // 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:<ol start="0"><li>`sycl::bundle_state::input`</li><li>`sycl::bundle_state::object`</li><li>`sycl::bundle_state::executable`</li></ol> |
+
+
+### [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:<ol start="0"><li>SPIR-V</li><li>PTX</li><li>AMDGCN</li></ol> |
+| "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. |
+