[SYCL][Docs] Add SYCLBIN feature and format design document

sycl/doc/design/SYCLBINDesign.asciidoc

@@ -0,0 +1,255 @@
+= SYCLBIN - A format for separately compiled SYCL device code
+
+This design document details the SYCLBIN binary format used for storing SYCL
+device binaries to be loaded dynamically by the SYCL runtime. It also details
+how the toolchain produces, links and packages these binaries, as well as how
+the SYCL runtime library handles files of this format.
+
+[[syclbin_format]]
+== SYCLBIN binary format
+
+The files produced by the new compilation path will follow the format described
+in this section. The intention of defining a new format for these is to give
+the DPC++ implementation an extendable and lightweight wrapper around the
+multiple modules and corresponding metadata captured in the SYCLBIN file.
+The content of the SYCLBIN may be contained as an entry in the offloading binary
+format produced by the clang-offload-packager, as described in
+link:../../clang/docs/ClangOffloadPackager.rst[ClangOffloadPackager.rst].
+
+The following illustration gives an overview of how the file format is
+structured.
+
+image::SYCLBIN_file_format_illustration.svg["SYCLBIN binary file format illustration", width=40%]
+
+=== Header
+
+The header segment appears as the first part of the SYCLBIN file. Like many
+other file-formats, it defines a magic number to help identify the format, which
+is 0x53594249 (or "SYBI".) Immediately following the magic number is the version
+number, which is used by SYCLBIN consumers when parsing data in the rest of the
+file.
+
+[cols="1,3,1"]
+|===
+| Type       | Description                                                        | Value variable
+
+| `uint32_t` | Magic number. (0x53594249)                                         | 
+| `uint32_t` | SYCLBIN version number.                                            |
+| `uint8_t`  | `sycl::bundle_state` corresponding to the contents of the SYCLBIN. |
+|===
+
+The `sycl::bundle_state` is an integer with the values as follows:
+
+[cols="3,1"]
+|===
+| `sycl::bundle_state` | Value
+
+| `input`              | 0
+| `object`             | 1
+| `executable`         | 2
+|===
+
+
+=== Body
+
+Immediately after the header is the body of the SYCLBIN file. The body consists
+of a list of abstract modules.
+
+[cols="1,3,1"]
+|===
+| Type       | Description                                | Value variable
+
+| `uint64_t` | Byte size of the list of abstract modules. | `B`
+| `B`        | List of abstract modules.                  |
+|===
+
+
+==== Abstract module
+
+Each abstract module represents a set of kernels, the corresponding metadata, 0
+or more IR modules containing these kernels, and 0 or more native device code
+images containing the kernels.
+
+[cols="1,3,1"]
+|===
+| Type       | Description                                     | Value variable
+
+| `uint64_t` | Byte size of the list of the metadata.          | `M`
+| `M`        | Module metadata.                                |
+| `uint64_t` | Byte size of list of IR modules.                | `IR`
+| `IR`       | List of IR modules.                             |
+| `uint64_t` | Byte size of list of native device code images. | `ND`
+| `ND`       | List of native device code images.              |
+|===
+
+
+===== Module metadata
+
+The module metadata contains the following information about the contents of the
+module.
+
+[cols="1,3,1"]
+|===
+| Type       | Description                                                    | Value variable
+
+| `uint32_t` | Byte size of the list of kernel names.                         | `K`
+| `K`        | List of kernel names. (String list)                            |
+| `uint32_t` | Byte size of the list of imported symbols.                     | `I`
+| `I`        | List of imported symbols. (String list)                        |
+| `uint32_t` | Byte size of the list of exported symbols.                     | `E`
+| `E`        | List of exported symbols. (String list)                        |
+| `uint32_t` | Byte size of property set data.                                | `P`
+| `P`        | Property set data.                                             |
+|===
+
+
+NOTE: Optional features used is embedded in the property set data.
+TODO: Consolidate and/or document the property set data in this document.
+
+====== String list
+
+A string list in this binary format consists of a `uint32_t` at the beginning
+containing the number of elements in the list, followed by that number of
+entries with the format:
+
+[cols="1,3,1"]
+|===
+| Type       | Description              | Value variable
+
+| `uint32_t` | Byte size of the string. | `S`
+| `S`        | String bytes.            |
+|===
+
+
+===== IR module
+
+An IR module contains the binary data for the corresponding module compiled to a
+given IR representation, identified by the IR type field.
+
+[cols="1,3,1"]
+|===
+| Type       | Description                    | Value variable
+
+| `uint8_t`  | IR type.                       |
+| `uint32_t` | Byte size of the raw IR bytes. | `IB`
+| `IB`       | Raw IR bytes.                  |
+|===
+
+TODO: Do we need a target-specific blob inside this structure? E.g. for CUDA we
+      may want to embed the SM version.
+
+
+====== IR types
+
+The IR types must be one of the following values:
+
+[cols="3,1"]
+|===
+| IR type | Value
+
+| SPIR-V  | 0
+|===
+
+
+===== Native device code image
+
+An native device code image contains the binary data for the corresponding
+module AOT compiled for a specific device, identified by the architecture
+string.
+
+[cols="1,3,1"]
+|===
+| Type       | Description                                      | Value variable
+
+| `uint32_t` | Byte size of the architecture string.            | `A`
+| `A`        | Architecture string.                             |
+| `uint32_t` | Byte size of the native device code image bytes. | `NB`
+| NB         | Native device code image bytes.                  |
+|===
+
+
+=== SYCLBIN version changelog
+
+The SYCLBIN format is subject to change, but any such changes must come with an
+increment to the version number in the header and a subsection to this section
+describing the change.
+
+==== Version 1
+
+ * Initial version of the layout.
+
+
+== Clang driver changes
+
+The clang driver needs to accept the following new flags:
+
+[cols="1,3"]
+|===
+| Option | Description
+
+| `-fsyclbin`
+| If this option is set, the output of the invocation is a SYCLBIN file with the
+  .syclbin file extension. This skips the host-compilation invocation of the
+  typical `-fsycl` pipeline, instead passing the output of the
+  clang-offloat-packager invocation to clang-linker-wrapper together with the
+  new `--syclbin` flag.
+
+  Setting this option will override `-fsycl` and `-fsycl-device-only`.
+
+  This option currently requires `--offload-new-driver` to be set.
+
+| `--offload-ir`
+| TODO
+
+| `--offload-rdc`
+| This is an alias of `-fgpu-rdc`.
+|===
+
+Additionally, `-fsycl-link` should work with .syclbin files. Semantics of how
+SYCLBIN files are linked together is yet to be specified.
+
+
+== clang-linker-wrapper changes
+
+The clang-linker-wrapper is responsible for doing post-processing and linking of
+device binaries, as described in
+link:OffloadDesign.rst[OffloadDesign.md].
+However, to support SYCLBIN files, the clang-linker-wrapper must be able to
+unpack an offload binary (as described in
+link:../../clang/docs/ClangOffloadPackager.rst[ClangOffloadPackager.rst])
+directly, instead of extracting it from a host binary. This should be done when
+a new flag, `--syclbin`, is passed. In this case, the clang-linker-wrapper is
+responsible to package the resulting device binaries and produced metadata into
+the format described in <<syclbin_format>>. Additionally, in this case the
+clang-linker-wrapper will skip the wrapping of the device code and the host code
+linking stage, as there is no host code to wrap the device code in and link.
+
+TODO: Describe the details of linking SYCLBIN files.
+
+
+== SYCL runtime library changes
+
+Using the interfaces from the
+link:../extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc[sycl_ext_oneapi_syclbin]
+extension, the runtime must be able to parse the SYCLBIN format, as described in
+<<syclbin_format>>. To avoid large amounts of code duplication, the runtime uses
+the implementation of SYCLBIN reading and writing implemented in LLVM.
+
+When creating a `kernel_bundle` from a SYCLBIN file, the runtime reads the
+contents of the SYCLBIN file and creates the corresponding data structure from
+it. In order for the SYCL runtime library's existing logic to use the binaries,
+the runtime then creates a collection of `sycl_device_binary_struct` objects and
+its constituents, pointing to the data in the parsed SYCLBIN object. Passing
+these objects to the runtime library's `ProgramManager` allows it to reuse the
+logic for compiling, linking and building SYCL binaries.
+
+In the other direction, users can request the "contents" of a `kernel_bundle`.
+When this is done, the runtime library must ensure that a SYCLBIN file is
+available for the contents of the `kernel_bundle` and must then write the
+SYCLBIN object to the corresponding binary representation in the format
+described in <<syclbin_format>>. In cases where the `kernel_bundle` was created
+with a SYCLBIN file, the SYCLBIN representation is immediately available and
+can be serialized directly. In other cases, the runtime library creates a new
+SYCLBIN object from the binaries associated with the `kernel_bundle`, then
+serializes it and returns the result.
+