[SYCL][UR][Docs] Add fast-link option for SYCLBIN

This commit adds the ability for doing "fast linking" of kernel bundles.
Fast linking lets the implementation use AOT binaries from the
underlying SYCLBIN files to dynamically link the images in the kernel
bundles.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

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<typename PropertyListT = empty_properties_t>
+kernel_bundle<bundle_state::executable>
+link(const std::vector<kernel_bundle<bundle_state::object>>& objectBundles,
+     const std::vector<device>& 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<typename PropertyListT = empty_properties_t>                        (1)
+kernel_bundle<bundle_state::executable>
+link(const kernel_bundle<bundle_state::object>& objectBundle,
+     const std::vector<device>& devs, PropertyListT props = {});
+
+template<typename PropertyListT = empty_properties_t>                        (2)
+kernel_bundle<bundle_state::executable>
+link(const std::vector<kernel_bundle<bundle_state::object>>& objectBundles,
+     PropertyListT props = {});
+
+template<typename PropertyListT = empty_properties_t>                        (3)
+kernel_bundle<bundle_state::executable>
+link(const kernel_bundle<bundle_state::object>& 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<fast_link_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.
+
+|====
+

sycl/include/sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/ext/oneapi/experimental/syclbin_properties.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/kernel_bundle.hpp>
 
@@ -24,6 +25,13 @@
 
 namespace sycl {
 inline namespace _V1 {
+
+namespace detail {
+__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
+link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
+          const std::vector<device> &Devs, bool FastLink);
+}
+
 namespace ext::oneapi::experimental {
 
 template <bundle_state State, typename PropertyListT = empty_properties_t>
@@ -77,6 +85,58 @@ get_kernel_bundle(const context &Ctxt, const std::filesystem::path &Filename,
 }
 #endif
 
+template <typename PropertyListT = empty_properties_t,
+          typename = std::enable_if_t<detail::all_are_properties_of_v<
+              sycl::detail::link_props, PropertyListT>>>
+kernel_bundle<bundle_state::executable>
+link(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
+     const std::vector<device> &Devs, PropertyListT Props = {}) {
+  std::vector<device> UniqueDevices =
+      sycl::detail::removeDuplicateDevices(Devs);
+
+  bool UseFastLink = [&]() {
+    if constexpr (Props.template has_property<fast_link>())
+      return Props.template get_property<fast_link>().value;
+    return false;
+  }();
+
+  sycl::detail::KernelBundleImplPtr Impl =
+      sycl::detail::link_impl(ObjectBundles, UniqueDevices, UseFastLink);
+  return detail::createSyclObjFromImpl<
+      kernel_bundle<sycl::bundle_state::executable>>(std::move(Impl));
+}
+
+template <typename PropertyListT = empty_properties_t,
+          typename = std::enable_if_t<detail::all_are_properties_of_v<
+              sycl::detail::link_props, PropertyListT>>>
+kernel_bundle<bundle_state::executable>
+link(const kernel_bundle<bundle_state::object> &ObjectBundle,
+     const std::vector<device> &Devs, PropertyListT Props = {}) {
+  return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
+              Devs, Props);
+}
+
+template <typename PropertyListT = empty_properties_t,
+          typename = std::enable_if_t<detail::all_are_properties_of_v<
+              sycl::detail::link_props, PropertyListT>>>
+kernel_bundle<bundle_state::executable>
+link(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
+     PropertyListT Props = {}) {
+  std::vector<sycl::device> IntersectDevices =
+      sycl::detail::find_device_intersection(ObjectBundles);
+  return link(ObjectBundles, IntersectDevices, Props);
+}
+
+template <typename PropertyListT = empty_properties_t,
+          typename = std::enable_if_t<detail::all_are_properties_of_v<
+              sycl::detail::link_props, PropertyListT>>>
+kernel_bundle<bundle_state::executable>
+link(const kernel_bundle<bundle_state::object> &ObjectBundle,
+     PropertyListT Props = {}) {
+  return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
+              ObjectBundle.get_devices(), Props);
+}
+
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/ext/oneapi/experimental/syclbin_properties.hpp

@@ -0,0 +1,39 @@
+//==-------- syclbin_properties.hpp - SYCLBIN and tooling properties -------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+
+namespace detail {
+struct link_props;
+} // namespace detail
+
+namespace ext::oneapi::experimental {
+
+/////////////////////////
+// PropertyT syclex::fast_link
+/////////////////////////
+struct fast_link
+    : detail::run_time_property_key<fast_link, detail::PropKind::FastLink> {
+  fast_link(bool DoFastLink = true) : value(DoFastLink) {}
+
+  bool value;
+};
+using fast_link_key = fast_link;
+
+template <>
+struct is_property_key_of<fast_link_key, sycl::detail::link_props>
+    : std::true_type {};
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -228,8 +228,9 @@ enum PropKind : uint32_t {
   InitialThreshold = 83,
   MaximumSize = 84,
   ZeroInit = 85,
+  FastLink = 86,
   // PropKindSize must always be the last value.
-  PropKindSize = 86,
+  PropKindSize = 87,
 };
 
 template <typename PropertyT> struct PropertyToKind {

sycl/include/sycl/sycl.hpp

@@ -133,6 +133,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/ext/oneapi/experimental/reduction_properties.hpp>
 #include <sycl/ext/oneapi/experimental/root_group.hpp>
 #include <sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp>
+#include <sycl/ext/oneapi/experimental/syclbin_properties.hpp>
 #include <sycl/ext/oneapi/experimental/tangle.hpp>
 #include <sycl/ext/oneapi/experimental/work_group_memory.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>

sycl/source/backend.cpp

@@ -231,7 +231,7 @@ make_kernel_bundle(ur_native_handle_t NativeHandle,
     case (UR_PROGRAM_BINARY_TYPE_NONE):
       if (State == bundle_state::object) {
         auto Res = Adapter.call_nocheck<UrApiKind::urProgramCompileExp>(
-            UrProgram, 1u, &Dev, nullptr);
+            UrProgram, 1u, &Dev, ur_exp_program_flags_t{}, nullptr);
         if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
           Res = Adapter.call_nocheck<UrApiKind::urProgramCompile>(
               ContextImpl.getHandleRef(), UrProgram, nullptr);
@@ -241,7 +241,7 @@ make_kernel_bundle(ur_native_handle_t NativeHandle,
 
       else if (State == bundle_state::executable) {
         auto Res = Adapter.call_nocheck<UrApiKind::urProgramBuildExp>(
-            UrProgram, 1u, &Dev, nullptr);
+            UrProgram, 1u, &Dev, ur_exp_program_flags_t{}, nullptr);
         if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
           Res = Adapter.call_nocheck<UrApiKind::urProgramBuild>(
               ContextImpl.getHandleRef(), UrProgram, nullptr);
@@ -261,8 +261,8 @@ make_kernel_bundle(ur_native_handle_t NativeHandle,
         Managed<ur_program_handle_t> UrLinkedProgram{Adapter};
         ur_program_handle_t ProgramsToLink[] = {UrProgram};
         auto Res = Adapter.call_nocheck<UrApiKind::urProgramLinkExp>(
-            ContextImpl.getHandleRef(), 1u, &Dev, 1u, ProgramsToLink, nullptr,
-            &UrLinkedProgram);
+            ContextImpl.getHandleRef(), 1u, &Dev, ur_exp_program_flags_t{}, 1u,
+            ProgramsToLink, nullptr, &UrLinkedProgram);
         if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
           Res = Adapter.call_nocheck<UrApiKind::urProgramLink>(
               ContextImpl.getHandleRef(), 1u, ProgramsToLink, nullptr,

sycl/source/detail/device_image_impl.hpp

@@ -761,7 +761,8 @@ class device_image_impl
 
     std::string XsFlags = extractXsFlags(BuildOptions, MRTCBinInfo->MLanguage);
     auto Res = Adapter.call_nocheck<UrApiKind::urProgramBuildExp>(
-        UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str());
+        UrProgram, DeviceVec.size(), DeviceVec.data(), ur_exp_program_flags_t{},
+        XsFlags.c_str());
     if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
       Res = Adapter.call_nocheck<UrApiKind::urProgramBuild>(
           ContextImpl.getHandleRef(), UrProgram, XsFlags.c_str());