[SYCL] Add fast-link option for SYCLBIN (#20174)

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.

Extension PR: #20271

---------

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,39 @@
 
 namespace sycl {
 inline namespace _V1 {
+
+namespace detail {
+__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
+link_impl(const kernel_bundle<bundle_state::object> *ObjectBundles,
+          size_t NumObjectBundles, const std::vector<device> &Devs,
+          bool FastLink);
+
+template <
+    typename PropertyListT = ext::oneapi::experimental::empty_properties_t,
+    typename = std::enable_if_t<
+        ext::oneapi::experimental::detail::all_are_properties_of_v<
+            ext::oneapi::experimental::detail::link_props, PropertyListT>>>
+kernel_bundle<bundle_state::executable>
+link_common(const kernel_bundle<bundle_state::object> *ObjectBundles,
+            size_t NumObjectBundles, const std::vector<device> &Devs,
+            PropertyListT Props = {}) {
+  std::vector<device> UniqueDevices = removeDuplicateDevices(Devs);
+
+  bool UseFastLink = [&]() {
+    if constexpr (Props.template has_property<
+                      ext::oneapi::experimental::fast_link>())
+      return Props.template get_property<ext::oneapi::experimental::fast_link>()
+          .value;
+    return false;
+  }();
+
+  KernelBundleImplPtr Impl =
+      link_impl(ObjectBundles, NumObjectBundles, UniqueDevices, UseFastLink);
+  return createSyclObjFromImpl<kernel_bundle<bundle_state::executable>>(
+      std::move(Impl));
+}
+} // namespace detail
+
 namespace ext::oneapi::experimental {
 
 template <bundle_state State, typename PropertyListT = empty_properties_t>
@@ -77,6 +111,46 @@ 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 = {}) {
+  return sycl::detail::link_common(ObjectBundles.data(), ObjectBundles.size(),
+                                   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 kernel_bundle<bundle_state::object> &ObjectBundle,
+     const std::vector<device> &Devs, PropertyListT Props = {}) {
+  return sycl::detail::link_common(&ObjectBundle, 1, 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 @@
+//==------------------------------------------------------------------------==//
+//
+// 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

@@ -138,6 +138,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/detail/helpers.hpp

@@ -108,6 +108,10 @@ template <typename iterator> class iterator_range {
   iterator_range(IterTy Begin, IterTy End, size_t Size)
       : Begin(Begin), End(End), Size(Size) {}
 
+  template <typename IterTy>
+  iterator_range(IterTy Begin, IterTy End)
+      : iterator_range(Begin, End, std::distance(Begin, End)) {}
+
   iterator_range() : iterator_range(iterator{}, iterator{}, 0) {}
 
   template <typename ContainerTy, typename = std::void_t<decltype(iterator{