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. + +|==== +