From 115eeaaef016f15b4a83ee5160395a5e0e8a1243 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 2 Oct 2025 07:29:04 -0700 Subject: [PATCH 1/4] [SYCL][Docs] Add link overload and fast-link property This commit adds a new link overload under the experimental namespace taking a compile-time property list. With this, a new property is added for enabling fast-linking, i.e. linking allowing AOT binaries in a SYCLBIN-based kernel bundle together with JIT binaries. These are all added as part of the sycl_ext_oneapi_syclbin extension. Signed-off-by: Larsen, Steffen --- .../proposed/sycl_ext_oneapi_syclbin.asciidoc | 139 ++++++++++++++++++ 1 file changed, 139 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc index 84d1ede6847a1..e308531453510 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -242,3 +242,142 @@ _{endnote}_] |==== +=== New free function for linking + +This extension adds the following new free functions to create and build a +kernel bundle in `ext_oneapi_source` state. + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +template +kernel_bundle +link(const std::vector>& objectBundles, + const std::vector& devs, PropertyListT props = {}); + +} // namespace sycl::ext::oneapi::experimental +---- +!==== + +_Constraints:_ Available only when `PropertyListT` is an instance of +`sycl::ext::oneapi::experimental::properties` which contains no properties +other than those listed below in the section "New properties for the +`link` function". + +_Effects:_ Duplicate device images from `objectBundles` are eliminated as though +they were joined via `join()`, then the remaining device images are translated +into one or more new device images of state `bundle_state::executable`, and a +new kernel bundle is created to contain these new device images. The new bundle +represents all of the kernels in `objectBundles` 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 linked and not represented +in the new bundle. + +The new bundle has the same associated context as those in `objectBundles`, and +the new bundle’s set of associated devices is `devs` (with duplicate devices +removed). + +_Returns:_ The new kernel bundle. + +_Throws:_ + +* An `exception` with the `errc::invalid` error code if the bundles in +`objectBundles` do not all have the same associated context. + +* An `exception` with the `errc::invalid` error code if any of the devices in +`devs` are not in the set of associated devices for any of the bundles in +`objectBundles` (as defined by `kernel_bundle::get_devices()`) or if the `devs` +vector is empty. + +* An `exception` with the `errc::build` error code if the online link operation +fails. + + +a| +[frame=all,grid=none] +!==== +a! +[source] +---- + +namespace sycl::ext::oneapi::experimental { + +template (1) +kernel_bundle +link(const kernel_bundle& objectBundle, + const std::vector& devs, PropertyListT props = {}); + +template (2) +kernel_bundle +link(const std::vector>& objectBundles, + PropertyListT props = {}); + +template (3) +kernel_bundle +link(const kernel_bundle& objectBundle, + PropertyListT props = {}); + +} // namespace sycl::ext::oneapi::experimental +---- +!==== + +_Effects (1):_ Equivalent to `link({objectBundle}, devs, props)`. + +_Effects (2):_ Equivalent to `link(objectBundles, devs, props)`, where `devs` is +the intersection of associated devices in common for all bundles in +`objectBundles`. + +_Effects (3):_ Equivalent to +`link({objectBundle}, objectBundle.get_devices(), props)`. + + +|==== + +=== New properties for the `link` function + +This extension adds the following properties, which can be used in conjunction +with the `link` function that is defined above: + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +struct fast_link { + fast_link(bool do_fast_link = true); (1) + + bool value; +}; +using fast_link_key = fast_link; + +template<> struct is_property_key : std::true_type {}; + +} // namespace sycl::ext::oneapi::experimental +---- +!==== + +This property instructs the `link` operation to do "fast linking". Enabling this +instructs the implementation to use device binary images that have been +pre-compiled. + +For example, SYCLBIN files may contain ahead-of-time compiled binary images +together with just-in-time compiled binary images, with the kernels and exported +functions potentially overlapping. When fast-linking is enabled, the +implementation will try to use the ahead-of-time compiled binary images over +their just-in-time compiled counterparts. + +_Effects (1):_ Creates a new `fast_link` property with a boolean value +indicating whether the `link` operation should do fast-linking. + +|==== + From 7173187901a9a31d5a3166703ba0d53e11ccf560 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 15 Oct 2025 04:29:37 -0700 Subject: [PATCH 2/4] Clarify trade-off Signed-off-by: Larsen, Steffen --- .../doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc index e308531453510..35cc4affd3652 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -370,6 +370,11 @@ This property instructs the `link` operation to do "fast linking". Enabling this instructs the implementation to use device binary images that have been pre-compiled. +Fast linking offers potentially faster link times, at the potential cost of +slower kernel execution time. The latter effect would mainly be due to link-time +optimizations the device compiler could have done during regular linking, that +cannot be done when doing fast linking. + For example, SYCLBIN files may contain ahead-of-time compiled binary images together with just-in-time compiled binary images, with the kernels and exported functions potentially overlapping. When fast-linking is enabled, the From a3beda8758a297b56d03dc83c6fa02983a497788 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 24 Oct 2025 04:29:07 -0700 Subject: [PATCH 3/4] Move new link function and property Signed-off-by: Larsen, Steffen --- .../sycl_ext_oneapi_kernel_compiler.asciidoc | 144 ++++++++++++++++++ .../proposed/sycl_ext_oneapi_syclbin.asciidoc | 144 ------------------ 2 files changed, 144 insertions(+), 144 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc index d21c5ff4c4394..2a20f107b19a3 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc @@ -618,6 +618,150 @@ property, but the duplicates have no effect. _{endnote}_] |==== +=== New free function for linking + +This extension adds the following new free functions to create and build a +kernel bundle in `ext_oneapi_source` state. + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +template +kernel_bundle +link(const std::vector>& objectBundles, + const std::vector& devs, PropertyListT props = {}); + +} // namespace sycl::ext::oneapi::experimental +---- +!==== + +_Constraints:_ Available only when `PropertyListT` is an instance of +`sycl::ext::oneapi::experimental::properties` which contains no properties +other than those listed below in the section "New properties for the +`link` function". + +_Effects:_ Duplicate device images from `objectBundles` are eliminated as though +they were joined via `join()`, then the remaining device images are translated +into one or more new device images of state `bundle_state::executable`, and a +new kernel bundle is created to contain these new device images. The new bundle +represents all of the kernels in `objectBundles` 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 linked and not represented +in the new bundle. + +The new bundle has the same associated context as those in `objectBundles`, and +the new bundle’s set of associated devices is `devs` (with duplicate devices +removed). + +_Returns:_ The new kernel bundle. + +_Throws:_ + +* An `exception` with the `errc::invalid` error code if the bundles in +`objectBundles` do not all have the same associated context. + +* An `exception` with the `errc::invalid` error code if any of the devices in +`devs` are not in the set of associated devices for any of the bundles in +`objectBundles` (as defined by `kernel_bundle::get_devices()`) or if the `devs` +vector is empty. + +* An `exception` with the `errc::build` error code if the online link operation +fails. + + +a| +[frame=all,grid=none] +!==== +a! +[source] +---- + +namespace sycl::ext::oneapi::experimental { + +template (1) +kernel_bundle +link(const kernel_bundle& objectBundle, + const std::vector& devs, PropertyListT props = {}); + +template (2) +kernel_bundle +link(const std::vector>& objectBundles, + PropertyListT props = {}); + +template (3) +kernel_bundle +link(const kernel_bundle& objectBundle, + PropertyListT props = {}); + +} // namespace sycl::ext::oneapi::experimental +---- +!==== + +_Effects (1):_ Equivalent to `link({objectBundle}, devs, props)`. + +_Effects (2):_ Equivalent to `link(objectBundles, devs, props)`, where `devs` is +the intersection of associated devices in common for all bundles in +`objectBundles`. + +_Effects (3):_ Equivalent to +`link({objectBundle}, objectBundle.get_devices(), props)`. + + +|==== + +=== New properties for the `link` function + +This extension adds the following properties, which can be used in conjunction +with the `link` function that is defined above: + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +struct fast_link { + fast_link(bool do_fast_link = true); (1) + + bool value; +}; +using fast_link_key = fast_link; + +template<> struct is_property_key : std::true_type {}; + +} // namespace sycl::ext::oneapi::experimental +---- +!==== + +This property instructs the `link` operation to do "fast linking". Enabling this +instructs the implementation to use device binary images that have been +pre-compiled. + +Fast linking offers potentially faster link times, at the potential cost of +slower kernel execution time. The latter effect would mainly be due to link-time +optimizations the device compiler could have done during regular linking, that +cannot be done when doing fast linking. + +For example, the binaries produced may contain ahead-of-time compiled binary +images together with just-in-time compiled binary images, with the kernels and +exported functions potentially overlapping. When fast-linking is enabled, the +implementation will try to use the ahead-of-time compiled binary images over +their just-in-time compiled counterparts. + +_Effects (1):_ Creates a new `fast_link` property with a boolean value +indicating whether the `link` operation should do fast-linking. + +|==== + === New constraint for kernel bundle member functions This extension adds the following constraint to some of the `kernel_bundle` diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc index 35cc4affd3652..84d1ede6847a1 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_syclbin.asciidoc @@ -242,147 +242,3 @@ _{endnote}_] |==== -=== New free function for linking - -This extension adds the following new free functions to create and build a -kernel bundle in `ext_oneapi_source` state. - -|==== -a| -[frame=all,grid=none] -!==== -a! -[source,c++] ----- -namespace sycl::ext::oneapi::experimental { - -template -kernel_bundle -link(const std::vector>& objectBundles, - const std::vector& devs, PropertyListT props = {}); - -} // namespace sycl::ext::oneapi::experimental ----- -!==== - -_Constraints:_ Available only when `PropertyListT` is an instance of -`sycl::ext::oneapi::experimental::properties` which contains no properties -other than those listed below in the section "New properties for the -`link` function". - -_Effects:_ Duplicate device images from `objectBundles` are eliminated as though -they were joined via `join()`, then the remaining device images are translated -into one or more new device images of state `bundle_state::executable`, and a -new kernel bundle is created to contain these new device images. The new bundle -represents all of the kernels in `objectBundles` 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 linked and not represented -in the new bundle. - -The new bundle has the same associated context as those in `objectBundles`, and -the new bundle’s set of associated devices is `devs` (with duplicate devices -removed). - -_Returns:_ The new kernel bundle. - -_Throws:_ - -* An `exception` with the `errc::invalid` error code if the bundles in -`objectBundles` do not all have the same associated context. - -* An `exception` with the `errc::invalid` error code if any of the devices in -`devs` are not in the set of associated devices for any of the bundles in -`objectBundles` (as defined by `kernel_bundle::get_devices()`) or if the `devs` -vector is empty. - -* An `exception` with the `errc::build` error code if the online link operation -fails. - - -a| -[frame=all,grid=none] -!==== -a! -[source] ----- - -namespace sycl::ext::oneapi::experimental { - -template (1) -kernel_bundle -link(const kernel_bundle& objectBundle, - const std::vector& devs, PropertyListT props = {}); - -template (2) -kernel_bundle -link(const std::vector>& objectBundles, - PropertyListT props = {}); - -template (3) -kernel_bundle -link(const kernel_bundle& objectBundle, - PropertyListT props = {}); - -} // namespace sycl::ext::oneapi::experimental ----- -!==== - -_Effects (1):_ Equivalent to `link({objectBundle}, devs, props)`. - -_Effects (2):_ Equivalent to `link(objectBundles, devs, props)`, where `devs` is -the intersection of associated devices in common for all bundles in -`objectBundles`. - -_Effects (3):_ Equivalent to -`link({objectBundle}, objectBundle.get_devices(), props)`. - - -|==== - -=== New properties for the `link` function - -This extension adds the following properties, which can be used in conjunction -with the `link` function that is defined above: - -|==== -a| -[frame=all,grid=none] -!==== -a! -[source,c++] ----- -namespace sycl::ext::oneapi::experimental { - -struct fast_link { - fast_link(bool do_fast_link = true); (1) - - bool value; -}; -using fast_link_key = fast_link; - -template<> struct is_property_key : std::true_type {}; - -} // namespace sycl::ext::oneapi::experimental ----- -!==== - -This property instructs the `link` operation to do "fast linking". Enabling this -instructs the implementation to use device binary images that have been -pre-compiled. - -Fast linking offers potentially faster link times, at the potential cost of -slower kernel execution time. The latter effect would mainly be due to link-time -optimizations the device compiler could have done during regular linking, that -cannot be done when doing fast linking. - -For example, SYCLBIN files may contain ahead-of-time compiled binary images -together with just-in-time compiled binary images, with the kernels and exported -functions potentially overlapping. When fast-linking is enabled, the -implementation will try to use the ahead-of-time compiled binary images over -their just-in-time compiled counterparts. - -_Effects (1):_ Creates a new `fast_link` property with a boolean value -indicating whether the `link` operation should do fast-linking. - -|==== - From 4ea81ae6203aba69b42614f308a5800ece55acdd Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Fri, 24 Oct 2025 15:12:00 +0200 Subject: [PATCH 4/4] Update sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc --- .../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc index 2a20f107b19a3..ec445627cb062 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc @@ -620,8 +620,9 @@ _{endnote}_] === New free function for linking -This extension adds the following new free functions to create and build a -kernel bundle in `ext_oneapi_source` state. +This extension adds the following new free functions to link kernel bundles in +`object` state. This differs from the regular `sycl::link` function the +additional property list argument. |==== a|