[SYCL][Doc] Add spec for "spirv" to kernel compiler (#11954)

Add an extension that defines support for online compiled kernels that
are written in binary SPIR-V.

Author: gmlueck

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -177,12 +177,26 @@ kernel_bundle<bundle_state::ext_oneapi_source> create_kernel_bundle_from_source(
   source_language lang,
   const std::string& source)
 
+kernel_bundle<bundle_state::ext_oneapi_source> create_kernel_bundle_from_source(
+  const context& ctxt,
+  source_language lang,
+  const std::vector<std::byte>& bytes)
+
 } // namespace sycl::ext::oneapi::experimental
 ----
 !====
 
-_Effects:_ Creates a new kernel bundle that represents a kernel that is defined
-by the source code string `source` using the language `lang`.
+_Preconditions:_ There are two overloads of this function: one that reads the
+source code of the kernel from an `std::string`, and one that reads the source
+code of the kernel from an `std::vector` of `std::byte`.
+Each source language `lang` specifies whether the language is text format or
+binary format, and the application must use the overload that corresponds to
+that format.
+
+_Effects:_ Creates a new kernel bundle that represents a kernel written in the
+source language `lang`, where the source code is contained either by `source`
+(if the source language is a text format) or by `bytes` (if the source language
+is binary format).
 The bundle is associated with the context `ctxt`, and kernels from this bundle
 may only be submitted to a queue that shares the same context.
 The bundle's set of associated devices is the set of devices contained in

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc

@@ -107,6 +107,13 @@ enum class source_language : /*unspecified*/ {
 } // namespace sycl::ext::oneapi::experimental
 ```
 
+=== Source code is text format
+
+Kernels written in the `opencl` language are text format.
+As a result, the application must use the overload of
+`create_kernel_bundle_from_source` taking `std::string` when creating a kernel
+bundle from this language.
+
 === Build options
 
 The `build_options` property accepts any of the compiler or linker options

sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc

@@ -0,0 +1,288 @@
+= sycl_ext_oneapi_kernel_compiler_spirv
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+:endnote: —{nbsp}end{nbsp}note
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2023-2023 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.
+OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos.
+
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 8 specification.
+All references below to the "core SYCL specification" or to section numbers in
+the SYCL specification refer to that revision.
+This extension also depends on the SPIR-V version 1.6, revision 2
+specification.
+All references below to the "SPIR-V specification" refer to that revision.
+
+This extension also depends on the following other SYCL extensions:
+
+* link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[
+  sycl_ext_oneapi_kernel_compiler]
+
+
+== Status
+
+This is a proposed extension specification, intended to gather community
+feedback.
+Interfaces defined in this specification may not be implemented yet or may be
+in a preliminary state.
+The specification itself may also change in incompatible ways before it is
+finalized.
+*Shipping software products should not rely on APIs defined in this
+specification.*
+
+
+== Overview
+
+This is an extension to
+link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[
+sycl_ext_oneapi_kernel_compiler], which allows an application to define a
+kernel as a SPIR-V binary module when dynamically compiling a kernel from
+source.
+One possible use case is an application that stores pre-compiled kernels as
+individual SPIR-V files.
+The application can load one of these files at runtime and then use this
+extension to enqueue the kernel and set its arguments.
+
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.
+An implementation supporting this extension must predefine the macro
+`SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV`
+to one of the values defined in the table below.
+Applications can test for the existence of this macro to determine if the
+implementation supports this feature, or applications can test the macro's
+value to determine which of the extension's features the implementation
+supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|The APIs of this experimental extension are not versioned, so the
+ feature-test macro always has this value.
+|===
+
+=== New source language enumerator
+
+This extension adds the `spirv` enumerator to the `source_language`
+enumeration, which indicates that a kernel bundle defines kernels as a SPIR-V
+binary module.
+
+```
+namespace sycl::ext::oneapi::experimental {
+
+enum class source_language : /*unspecified*/ {
+  // ...
+  spirv
+};
+
+} // namespace sycl::ext::oneapi::experimental
+```
+
+=== Source code is binary format
+
+Kernels written in the `spirv` language are SPIR-V binary modules.
+As a result, the application must use the overload of
+`create_kernel_bundle_from_source` taking `std::vector<std::byte>` when
+creating a kernel bundle from this language.
+
+=== Build options
+
+This extension does not specify any options that may be passed via the
+`build_options` property, however an implementation may allow
+implementation-defined options to be passed this way.
+
+=== SPIR-V execution environment
+
+The precise rules for interpreting a SPIR-V module are defined both by the
+SPIR-V specification and by the SPIR-V execution environment.
+For the purposes of this SYCL extension, the SPIR-V execution environment is
+defined by the SYCL backend.
+Each SYCL backend that supports this extension must provide a SPIR-V client API
+specification that formally defines the SPIR-V execution environment.
+
+[_Note:_ {dpcpp} provides two backends that can support this SYCL extension,
+and each has its own SPIR-V client API.
+For the Level Zero backend, see the
+https://spec.oneapi.io/level-zero/latest/core/SPIRV.html[SPIR-V Programming
+Guide].
+For the OpenCL backend, see the
+https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Env.html[
+OpenCL SPIR-V Environment Specification].
+_{endnote}_]
+
+=== Obtaining a kernel
+
+In SPIR-V, a kernel is represented as an *OpEntryPoint*, which has a literal
+_Name_ parameter.
+The `ext_oneapi_has_kernel` and `ext_oneapi_get_kernel` member functions
+identify a kernel using the name, exactly as it appears in the *OpEntryPoint*
+instruction.
+For example, if the kernel is defined this way in SPIR-V:
+
+```
+OpEntryPoint Kernel %20 "foo"
+```
+
+Then the application's host code can query for the kernel like so:
+
+```
+sycl::kernel_bundle<sycl::bundle_state::executable> kb = /*...*/;
+sycl::kernel k = kb.ext_oneapi_get_kernel("foo");
+```
+
+=== Passing kernel arguments
+
+The SPIR-V *OpEntryPoint* that defines a kernel has an associated *OpFunction*.
+This *OpFunction* is followed by a list of *OpFunctionParameter* instructions,
+one for each kernel argument.
+The following table defines the set of argument types that are supported by
+this extension, and explains how to pass each type of argument from SYCL.
+However, the set of supported argument types may be further limited by the
+backend's SPIR-V client API specification.
+
+[%header,cols="1,1"]
+|===
+|SPIR-V type
+|Corresponding SYCL type
+
+|*OpTypeInt*
+|A C++ type that is device copyable, which has the same width and data
+ representation.
+
+[_Note:_ Applications typically use a fixed-width integer type where the width
+matches the width of the *OpTypeInt*.
+_{endnote}_]
+
+|*OpTypeFloat*
+|A C++ type that is device copyable, which has the same width and data
+ representation.
+
+[_Note:_ Applications typically use `float` when the *OpTypeFloat* has a width
+of 32, `double` when the *OpTypeFloat* has a width of 64, and `sycl::half` when
+the *OpTypeFloat* has a width of 16.
+_{endnote}_]
+
+|*OpTypePointer* with storage class *CrossWorkgroup*
+|Either a pointer to USM memory or an `accessor` whose target is
+ `target::device`.
+
+|*OpTypePointer* with storage class *Workgroup*
+|A `local_accessor`.
+
+|*OpTypeStruct* whose member types are limited to *OpTypeInt*, *OpTypeFloat*,
+ and *OpTypePointer* (with storage class *CrossWorkgroup*).
+ The *OpTypeStruct* may contain members which are also *OpTypeStruct* so long
+ as its members are limited to the same types.
+|A C++ struct or class that is device copyable.
+ Each member variable must have the corresponding type as defined above,
+ except that an *OpTypePointer* member must correspond to a USM pointer.
+ It is not valid to pass an `accessor` for these members.
+|===
+
+When data allocated on the host is accessed by the kernel via a pointer, the
+application must ensure that the data has the same size and representation on
+the host and inside the SPIR-V module.
+
+[_Note:_ Applications should consider using the fixed-width integer types when
+allocating integer data that will be accessed by the kernel through a pointer
+because this helps ensure that the size of the integers on the host matches the
+size in the kernel.
+_{endnote}_]
+
+
+== Example
+
+The following example shows a simple SYCL program that loads a SPIR-V module
+from a file and then launches a kernel from that module.
+
+```
+#include <cstddef>
+#include <cstdint>
+#include <fstream>
+#include <vector>
+#include <sycl/sycl.hpp>
+
+namespace syclex = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue q;
+
+  // Read the SPIR-V module from disk.
+  std::ifstream spv_stream("my-kernel.spv", std::ios::binary);
+  spv_stream.seekg(0, std::ios::end);
+  size_t sz = spv_stream.tellg();
+  spv_stream.seekg(0);
+  std::vector<std::byte> spv(sz);
+  spv_stream.read((char*)spv.data(), sz);
+
+  // Create a kernel bundle from the binary SPIR-V.
+  sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
+    syclex::create_kernel_bundle_from_source(
+      q.get_context(),
+      syclex::source_language::spirv,
+      spv);
+
+  // Build the SPIR-V module for our device.
+  sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe =
+    syclex::build(kb_src);
+
+  // Get a "kernel" object representing the kernel from the SPIR-V module.
+  sycl::kernel k = kb_exe.ext_oneapi_get_kernel("my_kernel");
+
+  constexpr int N = 4;
+  int32_t input[N] = {0, 1, 2, 3};
+  int32_t output[N] = {};
+
+  sycl::buffer inputbuf(input, sycl::range{N});
+  sycl::buffer outputbuf(output, sycl::range{N});
+
+  q.submit([&](sycl::handler &cgh) {
+    sycl::accessor in{inputbuf, cgh, sycl::read_only};
+    sycl::accessor out{outputbuf, cgh, sycl::read_write};
+
+    // Set the values for the kernel arguments.
+    cgh.set_args(in, out);
+
+    // Invoke the kernel over a range.
+    cgh.parallel_for(sycl::range{N}, k);
+  });
+}
+```