From 1edfacf0ec50c0fecaa7211170475a0272de1ad9 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 27 Jan 2025 03:15:04 -0800 Subject: [PATCH 1/5] [SYCL][Docs] Add sycl_ext_oneapi_syclbin extension This commit adds the sycl_ext_oneapi_syclbin for loading a new SYCLBIN file format to `kernel_bundle`. Signed-off-by: Larsen, Steffen --- .../proposed/sycl_ext_oneapi_syclbin.asciidoc | 239 ++++++++++++++++++ 1 file changed, 239 insertions(+) create mode 100644 sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc new file mode 100644 index 0000000000000..c032559971a8d --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -0,0 +1,239 @@ += sycl_ext_oneapi_syclbin + +: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) 2025 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 9 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 following other SYCL extensions: + +* link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[ + sycl_ext_oneapi_kernel_compiler] +* link:../experimental/sycl_ext_oneapi_properties.asciidoc[ + sycl_ext_oneapi_properties] + + +== 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 extensions adds APIs, built upon the existing SYCL 2020 `kernel_bundle` +APIs, for loading precompiled "SYCLBIN" files. Using these, SYCL code can +dynamically load kernel binaries produced by the associated compiler, the format +and options for which are defined by the SYCL implementation. + + +== 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_SYCLBIN` 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. + + +=== New kernel_bundle interfaces + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source] +---- +namespace sycl::ext::oneapi::experimental { + +template +kernel_bundle get_kernel_bundle(const context& ctxt, + const std::vector& devs, + const sycl::span& bytes, + PropertyListT props = {}) + +#if __cplusplus >= 202002L +template +kernel_bundle get_kernel_bundle(const context& ctxt, + const std::vector& devs, + const std::span& bytes, + PropertyListT props = {}) (1) +#endif + +} +---- +!==== + +_Constraints:_ Available only when `State` is not `ext_oneapi_source`. +_Constraints (1):_ Available only when compiling for C++20 or later. + +_Effects:_ Creates a new kernel bundle containing the kernels from the SYCLBIN +data contained in `bytes` that are compatible with at least one of the devices +in `devs`. Any remaining kernels (those that are not compatible with any of the +devices in `devs`) are not represented in the new kernel bundle. + +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 and whose device +is in `devs`. + +_Returns:_ The newly created kernel bundle, which has `State` state. + +_Throws:_ + +* An `exception` with the `errc::invalid` error code if the contents of `bytes` + is not in the SYCLBIN format, as defined by the SYCL implementation. +* An `exception` with the `errc::invalid` error code if the `devs` vector is + empty. +* An `exception` with the `errc::invalid` error code if `State` is + `bundle_state::input` and any device in `ctxt.get_devices()` does not have + `aspect::online_compiler`. +* An `exception` with the `errc::invalid` error code if `State` is + `bundle_state::object` and any device in `ctxt.get_devices()` does not have + `aspect::online_linker`. +* An `exception` with the `errc::build` error code if `State` is + `bundle_state::object` or `bundle_state::executable`, if the implementation + needs to perform an online compile or link, and if the online compile or link + fails. + +a| +[frame=all,grid=none] +!==== +a! +[source] +---- +namespace sycl::ext::oneapi::experimental { + +template +kernel_bundle get_kernel_bundle(const context& ctxt, + const std::vector& devs, + const std::filesystem::path& filename, + PropertyListT props = {}) (1) + +template +kernel_bundle get_kernel_bundle(const context& ctxt, + const std::filesystem::path& filename, + PropertyListT props = {}) (2) + +} +---- +!==== + +_Effects (1):_ Creates a new kernel bundle containing the kernels inside the +SYCLBIN file located at `filename`. This is equivalent to +`get_kernel_bundle(ctxt, devs, data, props)` where `data` is the bytes read from +the SYCLBIN file at `filename`. + +_Effects (2)_: Equivalent to `get_kernel_bundle(ctxt, ctxt.get_devices(), filename, props)`. + +_Constraints:_ Available only when `State` is not `ext_oneapi_source`. + +_Returns:_ The newly created kernel bundle, which has `State` state. + +_Throws:_ + +* An `exception` with the `errc::invalid` error code if any of the devices in + `devs` is not one of devices contained by the context `ctxt` or is not a + descendent device of some device in `ctxt`. +* An `exception` with the `errc::invalid` error code if the `devs` vector is + empty. +* An `exception` with the `errc::invalid` error code if `State` is + `bundle_state::input` and any device in `ctxt.get_devices()` does not have + `aspect::online_compiler`. +* An `exception` with the `errc::invalid` error code if `State` is + `bundle_state::object` and any device in `ctxt.get_devices()` does not have + `aspect::online_linker`. +* An `exception` with the `errc::build` error code if `State` is + `bundle_state::object` or `bundle_state::executable`, if the implementation + needs to perform an online compile or link, and if the online compile or link + fails. + +|==== + +``` +namespace sycl { +template class kernel_bundle { +public: + ... + + std::vector ext_oneapi_get_content(); + + sycl::span ext_oneapi_get_content_sycl_view(); + +#if __cplusplus >= 202002L + std::span ext_oneapi_get_content_view(); +#endif + +}; +} +``` + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source] +---- +std::vector ext_oneapi_get_content() +sycl::span ext_oneapi_get_content_sycl_view() +std::span ext_oneapi_get_content_view() (1) +---- +!==== + +_Constraints:_ Available only when `State` is not `ext_oneapi_source`. +_Constraints (1):_ Available only when compiling for C++20 or later. + +_Returns:_ A container of bytes containing the data of the kernel bundle in the +SYCLBIN format. + +[_Note:_ The data returned by these member functions is not guaranteed to be the +same as was used when creating the `kernel_bundle` using the `get_kernel_bundle` +functions. +_{endnote}_] + +|==== + From 93befdb12685f9986d2a3f9e00d7a147cdae0db8 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 29 Jan 2025 05:57:28 -0800 Subject: [PATCH 2/5] Remove view getters Signed-off-by: Larsen, Steffen --- .../proposed/sycl_ext_oneapi_syclbin.asciidoc | 11 +---------- 1 file changed, 1 insertion(+), 10 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc index c032559971a8d..987ccf1dab3b5 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -201,12 +201,6 @@ public: std::vector ext_oneapi_get_content(); - sycl::span ext_oneapi_get_content_sycl_view(); - -#if __cplusplus >= 202002L - std::span ext_oneapi_get_content_view(); -#endif - }; } ``` @@ -219,15 +213,12 @@ a! [source] ---- std::vector ext_oneapi_get_content() -sycl::span ext_oneapi_get_content_sycl_view() -std::span ext_oneapi_get_content_view() (1) ---- !==== _Constraints:_ Available only when `State` is not `ext_oneapi_source`. -_Constraints (1):_ Available only when compiling for C++20 or later. -_Returns:_ A container of bytes containing the data of the kernel bundle in the +_Returns:_ A vector of bytes containing the data of the kernel bundle in the SYCLBIN format. [_Note:_ The data returned by these member functions is not guaranteed to be the From 7d42fe682b3a978fb9f8e8cc6521077d9451b62b Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 29 Jan 2025 22:34:35 -0800 Subject: [PATCH 3/5] Address comments Signed-off-by: Larsen, Steffen --- .../proposed/sycl_ext_oneapi_syclbin.asciidoc | 33 +++++++++++-------- 1 file changed, 20 insertions(+), 13 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc index 987ccf1dab3b5..9552d17f5c673 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -94,22 +94,20 @@ template kernel_bundle get_kernel_bundle(const context& ctxt, const std::vector& devs, const sycl::span& bytes, - PropertyListT props = {}) + PropertyListT props = {}); -#if __cplusplus >= 202002L +// Requires C++20 template kernel_bundle get_kernel_bundle(const context& ctxt, const std::vector& devs, const std::span& bytes, - PropertyListT props = {}) (1) -#endif + PropertyListT props = {}); } ---- !==== _Constraints:_ Available only when `State` is not `ext_oneapi_source`. -_Constraints (1):_ Available only when compiling for C++20 or later. _Effects:_ Creates a new kernel bundle containing the kernels from the SYCLBIN data contained in `bytes` that are compatible with at least one of the devices @@ -126,6 +124,8 @@ _Throws:_ * An `exception` with the `errc::invalid` error code if the contents of `bytes` is not in the SYCLBIN format, as defined by the SYCL implementation. +* An `exception` with the `errc::invalid` error code if the contents of `bytes` + is incompatible with `State`. * An `exception` with the `errc::invalid` error code if the `devs` vector is empty. * An `exception` with the `errc::invalid` error code if `State` is @@ -147,16 +147,16 @@ a! ---- namespace sycl::ext::oneapi::experimental { -template +template (1) kernel_bundle get_kernel_bundle(const context& ctxt, const std::vector& devs, const std::filesystem::path& filename, - PropertyListT props = {}) (1) + PropertyListT props = {}); -template +template (2) kernel_bundle get_kernel_bundle(const context& ctxt, const std::filesystem::path& filename, - PropertyListT props = {}) (2) + PropertyListT props = {}); } ---- @@ -175,6 +175,13 @@ _Returns:_ The newly created kernel bundle, which has `State` state. _Throws:_ +* A `std::ios_base::failure` exception if the function failed to access and read + the file specified by `filename`. +* An `exception` with the `errc::invalid` error code if the contents of the file + specified by `filename` is not in the SYCLBIN format, as defined by the SYCL + implementation. +* An `exception` with the `errc::invalid` error code if the contents of the file + specified by `filename` is incompatible with `State`. * An `exception` with the `errc::invalid` error code if any of the devices in `devs` is not one of devices contained by the context `ctxt` or is not a descendent device of some device in `ctxt`. @@ -219,11 +226,11 @@ std::vector ext_oneapi_get_content() _Constraints:_ Available only when `State` is not `ext_oneapi_source`. _Returns:_ A vector of bytes containing the data of the kernel bundle in the -SYCLBIN format. +SYCLBIN format for this implementation. -[_Note:_ The data returned by these member functions is not guaranteed to be the -same as was used when creating the `kernel_bundle` using the `get_kernel_bundle` -functions. +[_Note:_ If the `kernel_bundle` was created using the `get_kernel_bundle` from +a SYCLBIN file, the contents returned by this member function is not guaranteed +to be the same as the original SYCLBIN file. _{endnote}_] |==== From fef06a1156f912a362eba091416ac390e5bf3cac Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 30 Jan 2025 08:16:30 -0800 Subject: [PATCH 4/5] Small fixes and mention of state in SYCLBIN Signed-off-by: Larsen, Steffen --- .../proposed/sycl_ext_oneapi_syclbin.asciidoc | 21 +++++++++---------- 1 file changed, 10 insertions(+), 11 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc index 9552d17f5c673..f4b0ccfbf12fc 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -41,10 +41,8 @@ This extension is written against the SYCL 2020 revision 9 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 following other SYCL extensions: +This extension also depends on the following other SYCL extension: -* link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[ - sycl_ext_oneapi_kernel_compiler] * link:../experimental/sycl_ext_oneapi_properties.asciidoc[ sycl_ext_oneapi_properties] @@ -124,8 +122,8 @@ _Throws:_ * An `exception` with the `errc::invalid` error code if the contents of `bytes` is not in the SYCLBIN format, as defined by the SYCL implementation. -* An `exception` with the `errc::invalid` error code if the contents of `bytes` - is incompatible with `State`. +* An `exception` with the `errc::invalid` error code if the SYCLBIN read from + `bytes` is not in the `State` state. * An `exception` with the `errc::invalid` error code if the `devs` vector is empty. * An `exception` with the `errc::invalid` error code if `State` is @@ -162,6 +160,8 @@ kernel_bundle get_kernel_bundle(const context& ctxt, ---- !==== +_Constraints:_ Available only when `State` is not `ext_oneapi_source`. + _Effects (1):_ Creates a new kernel bundle containing the kernels inside the SYCLBIN file located at `filename`. This is equivalent to `get_kernel_bundle(ctxt, devs, data, props)` where `data` is the bytes read from @@ -169,8 +169,6 @@ the SYCLBIN file at `filename`. _Effects (2)_: Equivalent to `get_kernel_bundle(ctxt, ctxt.get_devices(), filename, props)`. -_Constraints:_ Available only when `State` is not `ext_oneapi_source`. - _Returns:_ The newly created kernel bundle, which has `State` state. _Throws:_ @@ -180,8 +178,8 @@ _Throws:_ * An `exception` with the `errc::invalid` error code if the contents of the file specified by `filename` is not in the SYCLBIN format, as defined by the SYCL implementation. -* An `exception` with the `errc::invalid` error code if the contents of the file - specified by `filename` is incompatible with `State`. +* An `exception` with the `errc::invalid` error code if the SYCLBIN read from + the file specified by `filename` is not in the `State` state. * An `exception` with the `errc::invalid` error code if any of the devices in `devs` is not one of devices contained by the context `ctxt` or is not a descendent device of some device in `ctxt`. @@ -226,10 +224,11 @@ std::vector ext_oneapi_get_content() _Constraints:_ Available only when `State` is not `ext_oneapi_source`. _Returns:_ A vector of bytes containing the data of the kernel bundle in the -SYCLBIN format for this implementation. +SYCLBIN format for this implementation. The corresponding SYCLBIN format will +be in `State` state. [_Note:_ If the `kernel_bundle` was created using the `get_kernel_bundle` from -a SYCLBIN file, the contents returned by this member function is not guaranteed +a SYCLBIN file, the contents returned by this member function are not guaranteed to be the same as the original SYCLBIN file. _{endnote}_] From c49694a84dcf193285b5af4f70c600a29ca8c911 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Sun, 2 Feb 2025 22:30:04 -0800 Subject: [PATCH 5/5] Add overview for state Signed-off-by: Larsen, Steffen --- .../extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc index f4b0ccfbf12fc..84d1ede6847a1 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -63,6 +63,14 @@ APIs, for loading precompiled "SYCLBIN" files. Using these, SYCL code can dynamically load kernel binaries produced by the associated compiler, the format and options for which are defined by the SYCL implementation. +Conversely, a SYCL implementation supporting this extension allows users to +create the binary contents in the SYCLBIN format from a `kernel_bundle` object, +even if that object was not created from a SYCLBIN file originally. As such, +despite the SYCL implementation defining the format of SYCLBIN files, the format +is guaranteed to contain the corresponding kernel bundle state of the SYCLBIN +contents, which must in turn match the state of any `kernel_bundle` object +created from it. + == Specification