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

steffenlarsen · 2025-09-23T14:04:12Z

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.

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]>

steffenlarsen · 2025-11-04T12:36:30Z

@intel/llvm-reviewers-runtime - This is ready for review! 🎉

steffenlarsen · 2025-11-06T15:51:22Z

@intel/llvm-reviewers-runtime - Friendly ping.

steffenlarsen · 2025-11-11T08:51:48Z

@intel/llvm-reviewers-runtime - Friendly ping.

vinser52 · 2025-11-11T18:35:47Z

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

+
+namespace detail {
+__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
+link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,


I do not like the idea of using the specific container (std::vector in that case) when we need to pass the range of objects to the ABI entry point.

Good point! It was copy-pasted from the existing link_impl without further thought, but I agree that passing ranges here is better. I decided to go with C-style pointer + size arguments to avoid sycl::span becoming a part of the ABI for this. It is converted to a span in the library however, to simplify the usage.

Actually, we already use sycl::span as the ABI entry-point parameter. Maybe we should use it here as well for consistency.

I'm not sure. We could, but I prefer C-style arrays to make the ABI breaking surface smaller once we have the actual std::span.

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

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

aelovikov-intel · 2025-11-12T16:26:19Z

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

+  return createSyclObjFromImpl<kernel_bundle<bundle_state::executable>>(
+      std::move(Impl));


Any particular reason to do this here instead of inside libsycl.so?

The function is templated, so we can't move all of link_common. If you mean change the return type of link_impl, it is currently a parallel (overload) to another link_impl, so it was chosen for consistency's sake.

you mean change the return type of link_impl

Yes, I think most of ABI interfaces should operate in terms of real SYCL objects.

aelovikov-intel · 2025-11-12T16:28:09Z

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

@@ -0,0 +1,39 @@
+//==-------- syclbin_properties.hpp - SYCLBIN and tooling properties -------==//


AFAIK, this part of the style guide has been dropped (having file name present in the very first line).

Do you mean it should be

Suggested change

//==-------- syclbin_properties.hpp - SYCLBIN and tooling properties -------==//

//==-------------------- SYCLBIN and tooling properties --------------------==//

or should the header be dropped entirely?

I think it's usually

//==-----------------------------------------------------------------------==// // // 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 // //===----------------------------------------------------------------------===// // /// \file <Description> ///

or something like that now.

sycl/source/detail/program_manager/program_manager.cpp

aelovikov-intel · 2025-11-12T18:06:32Z

sycl/source/detail/kernel_bundle_impl.hpp

+      // Collect all unique images.
+      std::vector<device_image_plain> DevImages;


Can you confirm that else part is just indentation change?

That is indeed the intention! 😄

aelovikov-intel · 2025-11-12T18:07:25Z

sycl/source/detail/kernel_bundle_impl.hpp

+
+    std::vector<const RTDeviceBinaryImage *> Result;
+    for (auto &SYCLBIN : MSYCLBINs) {
+      auto NativeBinImgs = SYCLBIN->getNativeBinaryImages(Dev);


Please use explicit type, it's not clear if there is a redundant copy or not.

It does indeed do some redundant copying. I've changed it to use move iterators here as well.

sycl/source/detail/syclbin.cpp

aelovikov-intel · 2025-11-12T18:08:55Z

sycl/test-e2e/SYCLBIN/link_object_fast_link.cpp

+//==-------------- link_input.cpp --- SYCLBIN extension tests --------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+


We don't do that in tests, AFAIK.

All the surrounding tests have it, but I don't mind either way. I can go remove it in the others after addressing the comments here.

#20626 removes the others.

sycl/test-e2e/SYCLBIN/link_object_fast_link.cpp

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

aelovikov-intel · 2025-11-13T15:33:47Z

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

+            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,


Do we need to throw an exception if size is zero?

aelovikov-intel · 2025-11-13T15:46:06Z

sycl/source/detail/kernel_bundle_impl.hpp

+  // Create a map between exported symbols and their indices in the device
+  // images collection.


Needs more clarification that LinkGraph expects the dependencies graph in terms of indices. Also being discussed offline on how to make it more readable.

One option is to add LinkGraph::addEdge, another is somethink like

// Format expected by the LinkGraph's ctor: std::vector<..> Dependencies; auto AddEdge = [&Dependencies](...) { ... };

aelovikov-intel · 2025-11-13T15:48:38Z

sycl/source/detail/kernel_bundle_impl.hpp

+      // When doing fast-linking, we insert the suitable AOT binaries from the
+      // object bundles. This needs to be done per-device, as AOT binaries may
+      // not be compatible across different architectures.


Don't know the detail, but any chance the comment could be lifter to right before if (FastLink) to describe similarities/differences between two paths?

aelovikov-intel · 2025-11-13T15:49:49Z

sycl/source/detail/kernel_bundle_impl.hpp

+          const std::shared_ptr<detail::kernel_bundle_impl> &ObjectBundleImpl =
+              getSyclObjImpl(ObjectBundle);


Suggested change

const std::shared_ptr<detail::kernel_bundle_impl> &ObjectBundleImpl =

getSyclObjImpl(ObjectBundle);

detail::kernel_bundle_impl>&ObjectBundleImpl = *getSyclObjImpl(ObjectBundle);

aelovikov-intel · 2025-11-13T15:51:59Z

sycl/source/detail/kernel_bundle_impl.hpp

+            // If any of the exported symbols overlap with an AOT binary, skip
+            // this image.


Maybe add extra comment describing when that may happen and why? Same below.

aelovikov-intel · 2025-11-13T16:00:25Z

sycl/source/detail/kernel_bundle_impl.hpp

+      sycl::span<const device_image_plain> AllDevImgsSpan(AllDevImgs.data(),
+                                                          AllDevImgs.size());


Why do you need this? span is constructible from std::vector: https://godbolt.org/z/bxGde7WPr, or rather from Container,

llvm/sycl/include/sycl/sycl_span.hpp

Lines 464 to 469 in 20dd0bc

template <class _Container>

_SYCL_SPAN_INLINE_VISIBILITY constexpr span(

_Container &__c,

std::enable_if_t<__is_span_compatible_container<_Container, _Tp>::value,

std::nullptr_t> = nullptr)

: __data{std::data(__c)}, __size{(size_type)std::size(__c)} {}

aelovikov-intel · 2025-11-13T16:03:34Z

sycl/source/detail/kernel_bundle_impl.hpp

-      const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
-      devices_range Devs, const property_list &PropList, private_tag)
+  kernel_bundle_impl(sycl::span<const kernel_bundle<bundle_state::object>,
+                                sycl::dynamic_extent> &ObjectBundles,


Drop sycl::dynamic_extent.

aelovikov-intel · 2025-11-13T16:03:42Z

sycl/source/kernel_bundle.cpp

 link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
          const std::vector<device> &Devs, const property_list &PropList) {
-  return detail::kernel_bundle_impl::create(ObjectBundles, Devs, PropList);
+  sycl::span<const kernel_bundle<bundle_state::object>, sycl::dynamic_extent>


Drop sycl::dynamic_extent.

aelovikov-intel · 2025-11-13T16:03:48Z

sycl/source/kernel_bundle.cpp

+link_impl(const kernel_bundle<bundle_state::object> *ObjectBundles,
+          size_t NumObjectBundles, const std::vector<device> &Devs,
+          bool FastLink) {
+  sycl::span<const kernel_bundle<bundle_state::object>, sycl::dynamic_extent>


Drop sycl::dynamic_extent.

steffenlarsen had a problem deploying to WindowsCILock September 23, 2025 14:04 — with GitHub Actions Error

steffenlarsen force-pushed the steffen/dynamic_linking branch from 84ffa78 to b8c44d9 Compare September 23, 2025 14:34

steffenlarsen had a problem deploying to WindowsCILock September 23, 2025 14:35 — with GitHub Actions Error

steffenlarsen force-pushed the steffen/dynamic_linking branch from b8c44d9 to cc05fe9 Compare September 23, 2025 15:05

steffenlarsen temporarily deployed to WindowsCILock September 23, 2025 15:05 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock September 23, 2025 16:06 — with GitHub Actions Inactive

steffenlarsen mentioned this pull request Sep 25, 2025

[SYCL][SYCLBIN] Select best images based on state #20197

Merged

steffenlarsen force-pushed the steffen/dynamic_linking branch from cc05fe9 to aa08b13 Compare September 29, 2025 16:08

steffenlarsen temporarily deployed to WindowsCILock September 29, 2025 16:08 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock September 29, 2025 17:01 — with GitHub Actions Inactive

steffenlarsen force-pushed the steffen/dynamic_linking branch from aa08b13 to f84e78e Compare October 28, 2025 08:40

steffenlarsen temporarily deployed to WindowsCILock October 28, 2025 08:40 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock October 28, 2025 09:13 — with GitHub Actions Inactive

steffenlarsen force-pushed the steffen/dynamic_linking branch from f84e78e to 6fe8e6b Compare October 28, 2025 10:45

steffenlarsen temporarily deployed to WindowsCILock October 28, 2025 10:46 — with GitHub Actions Inactive

steffenlarsen changed the title ~~[SYCL][UR][Docs] Add fast-link option for SYCLBIN~~ [SYCL] Add fast-link option for SYCLBIN Oct 28, 2025

steffenlarsen temporarily deployed to WindowsCILock October 28, 2025 11:17 — with GitHub Actions Inactive

steffenlarsen marked this pull request as ready for review October 29, 2025 10:38

steffenlarsen requested a review from a team as a code owner October 29, 2025 10:38

steffenlarsen requested a review from vinser52 October 29, 2025 10:38

vinser52 reviewed Nov 11, 2025

View reviewed changes

sycl/include/sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp Outdated Show resolved Hide resolved

steffenlarsen added 2 commits November 11, 2025 22:53

Merge remote-tracking branch 'intel/sycl' into steffen/dynamic_linking

7111c47

Switch to passing pointer to bundles

30017be

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

steffenlarsen had a problem deploying to WindowsCILock November 12, 2025 07:24 — with GitHub Actions Failure

steffenlarsen temporarily deployed to WindowsCILock November 12, 2025 07:50 — with GitHub Actions Inactive

steffenlarsen had a problem deploying to WindowsCILock November 12, 2025 07:50 — with GitHub Actions Failure

steffenlarsen temporarily deployed to WindowsCILock November 12, 2025 07:50 — with GitHub Actions Inactive

Fix Windows symbol

cc75ee8

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

steffenlarsen had a problem deploying to WindowsCILock November 12, 2025 08:21 — with GitHub Actions Failure

steffenlarsen temporarily deployed to WindowsCILock November 12, 2025 08:54 — with GitHub Actions Inactive

steffenlarsen had a problem deploying to WindowsCILock November 12, 2025 08:54 — with GitHub Actions Failure

Add back unintentionally removed newline

0d62963

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

steffenlarsen temporarily deployed to WindowsCILock November 12, 2025 09:51 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock November 12, 2025 10:23 — with GitHub Actions Inactive

vinser52 requested a review from aelovikov-intel November 12, 2025 15:15

aelovikov-intel reviewed Nov 12, 2025

View reviewed changes

steffenlarsen added 2 commits November 12, 2025 23:57

Address comments

1400a6d

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

Merge remote-tracking branch 'intel/sycl' into steffen/dynamic_linking

81987ad

steffenlarsen had a problem deploying to WindowsCILock November 13, 2025 07:58 — with GitHub Actions Error

Move from native bins collection

638aa48

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

steffenlarsen temporarily deployed to WindowsCILock November 13, 2025 08:12 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock November 13, 2025 08:45 — with GitHub Actions Inactive

steffenlarsen mentioned this pull request Nov 13, 2025

[SYCL][E2E][NFC] Drop copyright headers in SYCLBIN E2E tests #20626

Merged

aelovikov-intel reviewed Nov 13, 2025

View reviewed changes

		return createSyclObjFromImpl<kernel_bundle<bundle_state::executable>>(
		std::move(Impl));

		@@ -0,0 +1,39 @@
		//==-------- syclbin_properties.hpp - SYCLBIN and tooling properties -------==//

	//==-------- syclbin_properties.hpp - SYCLBIN and tooling properties -------==//
	//==-------------------- SYCLBIN and tooling properties --------------------==//

		// Collect all unique images.
		std::vector<device_image_plain> DevImages;

		// Create a map between exported symbols and their indices in the device
		// images collection.

		const std::shared_ptr<detail::kernel_bundle_impl> &ObjectBundleImpl =
		getSyclObjImpl(ObjectBundle);

	const std::shared_ptr<detail::kernel_bundle_impl> &ObjectBundleImpl =
	getSyclObjImpl(ObjectBundle);
	detail::kernel_bundle_impl>&ObjectBundleImpl = *getSyclObjImpl(ObjectBundle);

		// If any of the exported symbols overlap with an AOT binary, skip
		// this image.

		sycl::span<const device_image_plain> AllDevImgsSpan(AllDevImgs.data(),
		AllDevImgs.size());

	template <class _Container>
	_SYCL_SPAN_INLINE_VISIBILITY constexpr span(
	_Container &__c,
	std::enable_if_t<__is_span_compatible_container<_Container, _Tp>::value,
	std::nullptr_t> = nullptr)
	: __data{std::data(__c)}, __size{(size_type)std::size(__c)} {}

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

Are you sure you want to change the base?

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

Uh oh!

Conversation

steffenlarsen commented Sep 23, 2025

Uh oh!

steffenlarsen commented Nov 4, 2025

Uh oh!

steffenlarsen commented Nov 6, 2025

Uh oh!

steffenlarsen commented Nov 11, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants