[SYCL][Doc] Add spec to get device image backend content

sycl/doc/extensions/proposed/sycl_ext_oneapi_device_image_content.asciidoc

@@ -0,0 +1,254 @@
+= sycl_ext_oneapi_device_image_content
+
+: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) 2024 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.
+
+
+== 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 extension adds a mechanism to obtain the raw content of the device images
+that are in a kernel bundle.
+The format of this content is implementation-defined, so applications that make
+use of this extension are normally not portable to other implementations of
+SYCL.
+
+
+== 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_DEVICE_IMAGE_CONTENT` 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 member functions in the `device_image` class
+
+This extension adds the following member functions to the `device_image` class.
+
+[source,c++]
+----
+namespace sycl {
+
+template <bundle_state State>
+class device_image {
+ public:
+  std::vector<std::byte> ext_oneapi_get_content() const;
+  std::span<std::byte> ext_oneapi_get_content_view() const;
+
+  /*...*/
+};
+
+} // namespace sycl
+----
+
+'''
+
+[frame=all,grid=none,separator="@"]
+!====
+a@
+[source,c++]
+----
+std::vector<std::byte> ext_oneapi_get_content() const;
+----
+!====
+
+_Returns:_ A copy of the raw bytes for this device image.
+The format of this data is implementation-defined.
+See below for a description of this format on {dpcpp}.
+
+'''
+
+[frame=all,grid=none,separator="@"]
+!====
+a@
+[source,c++]
+----
+std::span<std::byte> ext_oneapi_get_content_view() const;
+----
+!====
+
+Available only when the compiler is {cpp}20 or higher.
+
+_Returns:_ A view of the raw bytes for this device image.
+The data behind this view has the same lifetime as the `kernel_bundle` that
+contains this `device_image`.
+
+'''
+
+
+== Device image format for {dpcpp}
+
+This section is non-normative and applies only to the {dpcpp} implementation.
+The format of the data returned by `device_image::ext_oneapi_get_content` and
+`device_image::ext_oneapi_get_content_view` depends on the backend of the kernel
+bundle that contains the device image and also on the `State` of the device
+image.
+
+=== Format on Level Zero
+
+The format depends on the `State` of the device image:
+
+* `bundle_state::input`: The format could change in the future, but it is
+  currently a SPIR-V module representing one or more kernels.
+  This SPIR-V module may be partially linked, with references to internal
+  library functions that are not defined in the SPIR-V module.
+
+* `bundle_state::object`: The same format as `bundle_state::input`.
+
+* `bundle_state::executable`: The device image content is native ISA for the
+  device, which can be passed to `zeModuleCreate` as `ZE_MODULE_FORMAT_NATIVE`
+  format.
+
+:ref1: ../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc#level-zero-and-opencl-compatibility
+
+[_Note:_ The interface to kernels in the device image is not defined in the
+general case, which means there is no portable way to invoke kernels from a
+Level Zero module that is created from the raw device image content.
+However, see link:{ref1}[here] for a limited case where this portability is
+guaranteed.
+_{endnote}_]
+
+=== Format on OpenCL
+
+The format depends on the `State` of the device image:
+
+* `bundle_state::input`: The format could change in the future, but it is
+  currently a SPIR-V module as described above for Level Zero.
+
+* `bundle_state::object`: The device image content is an unspecified format that
+  is created by calling `clCompileProgram`.
+
+* `bundle_state::executable`: The device image content is executable binary
+  device code representing one or more kernels, which can be passed to
+  `clCreateProgramWithBinary`.
+
+[_Note:_ The interface to kernels in the device image is not defined in the
+general case, which means there is no portable way to invoke kernels from a
+OpenCL `cl_program` object that is created from the raw device image content.
+However, see link:{ref1}[here] for a limited case where this portability is
+guaranteed.
+_{endnote}_]
+
+=== Format on CUDA
+
+The format depends on the `State` of the device image:
+
+* `bundle_state::input`: The device image content is a PTX module representing
+  one or more kernels.
+
+* `bundle_state::object`: ???
+
+* `bundle_state::executable`: The device image content is a CUBIN module
+  representing one or more kernels.
+
+
+== Example
+
+:ref2: ../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc
+
+A kernel bundle can contain multiple device images with different
+representations of the same kernel for different devices.
+This example shows how to get the device image content for a particular kernel
+for a particular device.
+Note that this example also uses the kernel syntax described in link:{ref2}[
+sycl_ext_oneapi_free_function_kernels], but it is not necessary to define
+kernels in that syntax when using this extension.
+
+[source,c++]
+----
+#include <sycl/sycl.hpp>
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void iota(float start, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
+  ptr[id] = start + static_cast<float>(id);
+}
+
+void main() {
+  sycl::device d;
+  sycl::queue q{d};
+  sycl::context ctxt = q.get_context();
+
+  // Get a kernel bundle that contains the kernel "iota".
+  sycl::kernel_id iota = syclexp::get_kernel_id<iota>();
+  auto exe_bndl =
+    sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctxt, {iota});
+
+  std::vector<std::byte> bytes;
+  for (auto& img: bundle) {
+    // Search for the device image that contains "iota" for this device.
+    if (img.has_kernel(iota, dev)) {
+      bytes = img.ext_oneapi_get_content();
+      break;
+    }
+  }
+}
+----

sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc

@@ -773,15 +773,14 @@ int main() {
 ```
 
 
+[[level-zero-and-opencl-compatibility]]
 == {dpcpp} guaranteed compatibility with Level Zero and OpenCL backends
 
 The contents of this section are non-normative and apply only to the {dpcpp}
 implementation.
 Kernels written using the free function kernel syntax can be submitted to a
 device by using the Level Zero or OpenCL backends, without going through the
 SYCL host runtime APIs.
-This works only when the kernel is AOT compiled to native device code using the
-`-fsycl-targets` compiler option.
 
 The interface to the kernel in the native device code module is only guaranteed
 when the kernel adheres to the following restrictions:
@@ -795,6 +794,17 @@ when the kernel adheres to the following restrictions:
 * The translation unit containing the kernel is compiled with the
   `-fno-sycl-dead-args-optimization` option.
 
+In order to invoke a kernel using Level Zero or OpenCL, the application must
+first obtain the raw native ISA content of the device image that contains the
+kernel.
+One way to do this is by using
+link:../proposed/sycl_ext_oneapi_device_image_content.asciidoc[
+sycl_ext_oneapi_device_image_content] to obtain the content of the device image
+in `bundle_state::executable` state.
+It is also possible to compile the application in AOT mode via the
+`-fsycl-targets` compiler option and then extract the device image from the
+executable file.
+
 Both Level Zero and OpenCL identify a kernel via a _name_ string.
 (See `zeKernelCreate` and `clCreateKernel` in their respective specifications.)
 When a kernel is defined according to the restrictions above, the _name_ is