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

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

@@ -134,6 +134,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());