[SYCL] 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/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

@@ -137,6 +137,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>