[SYCL] Implement loading SYCLBIN into kernel_bundle #18949

steffenlarsen · 2025-06-12T09:26:04Z

This commit implements the functionality for loading SYCLBIN files into kernel bundles. This is done by mimicking the structure of regular device binaries, then letting the existing functionality handle compiling and linking.

This implements part of the sycl_ext_oneapi_syclbin extension.

Note that parts of this implementation uses functionality copied from LLVMSupport and LLVMObject. Eventually they should be replaced in favor of using the LLVM libraries directly.

This commit implements the functionality for loading SYCLBIN files into kernel bundles. This is done by mimicing the structure of regular device binaries, then letting the existing functionality handle compiling and linking. This implements part of the sycl_ext_oneapi_syclbin extension. Signed-off-by: Larsen, Steffen <[email protected]>

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

sycl/source/detail/device_binary_image.cpp

cperkinsintel · 2025-06-23T19:48:54Z

sycl/source/detail/kernel_bundle_impl.hpp

@@ -463,12 +464,27 @@ class kernel_bundle_impl
            "Not all input bundles have the same set of associated devices.");
    }

+    // Pre-count and reserve space in vectors.


question: Is constant time access to all three of these really important?

If it is, let's continue with std::vector. But if it isn't, maybe we should just use std::list? With std::list we don't need to "reserve", we get constant time insertion and deletion. Deleting/erase an element doesn't cause move on the other elements. There are a lot of advantages. Of course, lookup, even by index, becomes linear time instead of constant time.

They are occasionally all iterated over, which is faster with std::vector. std::list would indeed have constant selective deletion, insertion and preservation of memory, but these collections should not change after creation. The merging is only done when joining/linking kernel bundles, as the results are the combination of all the kernel bundles.

dm-vodopyanov

Looks awesome! 🎉

Just some non-blocking comments which can be addressed in a separate PR

sycl/source/detail/device_binary_image.cpp

sycl/source/detail/syclbin.cpp

dm-vodopyanov · 2025-06-24T16:02:28Z

sycl/source/detail/syclbin.cpp

+      reinterpret_cast<const OffloadBinaryEntryType *>(Data +
+                                                       Header->EntryOffset);
+
+  if (Entry->ImageKind != /*IMG_SYCLBIN*/ 6)


Should we move this constant somewhere? Where did 6 come from?

It's sadly a little bit of a hack, stemming from https://github.com/intel/llvm/blob/sycl/llvm/include/llvm/Object/OffloadBinary.h#L50. Eventually this file will disappear in favor of the LLVMObject version of SYCLBIN parsing, but for now we can't access that enum. Do you think it would make sense to make a constexpr uint16_t IMG_SYCLBIN = 6; and use that? I'm not sure how much it adds as it's still just a magic value, but maybe it's easier to make the intention clear. Alternatively, I can add a comment.

Do you think this value will be used somewhere else in this temporary file? If not, I guess this is fine to keep as is, otherwise constexpr variable is good to be added, but overall this is nit

We probably want to use it when we repack eventually too, but I hope we can have these temps gone by that time.

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

dm-vodopyanov · 2025-06-24T16:17:34Z

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

+  {
+    std::ifstream FileStream{Filename, std::ios::binary};
+    if (!FileStream.is_open())
+      throw sycl::exception(make_error_code(errc::invalid),


From the spec:

_Throws:_ * A `std::ios_base::failure` exception if the function failed to access and read the file specified by `filename`.

Ah, good point!

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

…bundle_impl` Follow-up for intel#18899. Also adds proper `private_tag` argument for the ctor added in intel#18949 that missed that.

…bundle_impl` Follow-up for intel#18899. Also adds proper `private_tag` argument for the ctor added in intel#18949 that missed that. Also pass `span` by value while on it.

…bundle_impl` (#19185) Follow-up for #18899. Also adds proper `private_tag` argument for the ctor added in #18949 that missed that. Also pass `span` by value while on it.

steffenlarsen requested a review from a team as a code owner June 12, 2025 09:26

steffenlarsen requested a review from uditagarwal97 June 12, 2025 09:26

steffenlarsen had a problem deploying to WindowsCILock June 12, 2025 09:26 — with GitHub Actions Error

steffenlarsen requested review from jopperm, sommerlukas and uditagarwal97 and removed request for uditagarwal97 June 12, 2025 09:26

Fix formatting

ec21a10

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

steffenlarsen had a problem deploying to WindowsCILock June 12, 2025 09:28 — with GitHub Actions Failure

Address formatting and warnings

7061d2c

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

steffenlarsen had a problem deploying to WindowsCILock June 12, 2025 11:23 — with GitHub Actions Failure

Even more pedantic formatting and errors

9bb13aa

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

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

steffenlarsen added 2 commits June 12, 2025 06:49

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

c9f17c2

Rebase and fix warning

258ecee

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

steffenlarsen requested a review from a team as a code owner June 12, 2025 13:50

steffenlarsen had a problem deploying to WindowsCILock June 12, 2025 13:51 — with GitHub Actions Error

Fix the right file

4700308

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

steffenlarsen had a problem deploying to WindowsCILock June 12, 2025 14:00 — with GitHub Actions Failure

steffenlarsen added 2 commits June 12, 2025 07:40

Fix windows build failure

2130901

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

Revert tooling changes

c49ddf4

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

steffenlarsen had a problem deploying to WindowsCILock June 12, 2025 14:47 — with GitHub Actions Error

steffenlarsen removed the request for review from a team June 12, 2025 14:47

Exclude CUDA and HIP for now

af8e38e

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

steffenlarsen had a problem deploying to WindowsCILock June 12, 2025 15:30 — with GitHub Actions Failure

steffenlarsen temporarily deployed to WindowsCILock June 12, 2025 17:55 — with GitHub Actions Inactive

Avoid charconv for RHEL builds

9b632ba

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

steffenlarsen had a problem deploying to WindowsCILock June 13, 2025 08:06 — with GitHub Actions Failure

steffenlarsen had a problem deploying to WindowsCILock June 18, 2025 07:10 — with GitHub Actions Failure

steffenlarsen had a problem deploying to WindowsCILock June 18, 2025 07:37 — with GitHub Actions Failure

Revert lazy image size change

abc3a2b

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

steffenlarsen had a problem deploying to WindowsCILock June 18, 2025 09:05 — with GitHub Actions Failure

steffenlarsen commented Jun 18, 2025

View reviewed changes

sycl/source/detail/device_binary_image.cpp Show resolved Hide resolved

steffenlarsen had a problem deploying to WindowsCILock June 18, 2025 09:30 — with GitHub Actions Failure

jopperm approved these changes Jun 18, 2025

View reviewed changes

steffenlarsen mentioned this pull request Jun 19, 2025

[SYCL] Fix unloading of shared images on Windows #19074

Closed

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

f333fbe

steffenlarsen had a problem deploying to WindowsCILock June 20, 2025 10:29 — with GitHub Actions Failure

steffenlarsen had a problem deploying to WindowsCILock June 20, 2025 10:54 — with GitHub Actions Failure

steffenlarsen requested a review from cperkinsintel June 23, 2025 11:56

cperkinsintel reviewed Jun 23, 2025

View reviewed changes

dm-vodopyanov approved these changes Jun 24, 2025

View reviewed changes

steffenlarsen added 2 commits June 24, 2025 22:29

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

e9dd9a9

Address comments

b248a6f

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

steffenlarsen temporarily deployed to WindowsCILock June 25, 2025 05:58 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock June 25, 2025 06:23 — with GitHub Actions Inactive

steffenlarsen requested a review from cperkinsintel June 25, 2025 11:44

cperkinsintel approved these changes Jun 25, 2025

View reviewed changes

steffenlarsen merged commit 0ab45d9 into intel:sycl Jun 26, 2025
25 checks passed

steffenlarsen deleted the steffen/load_syclbin_kb branch June 26, 2025 05:09

aelovikov-intel mentioned this pull request Jun 26, 2025

[NFC][SYCL] More cleanup after enable_shared_from_this for kernel_bundle_impl #19185

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[SYCL] Implement loading SYCLBIN into kernel_bundle #18949

[SYCL] Implement loading SYCLBIN into kernel_bundle #18949

Uh oh!

steffenlarsen commented Jun 12, 2025

Uh oh!

Uh oh!

cperkinsintel Jun 23, 2025

Uh oh!

steffenlarsen Jun 24, 2025

Uh oh!

dm-vodopyanov left a comment

Uh oh!

Uh oh!

Uh oh!

dm-vodopyanov Jun 24, 2025

Uh oh!

steffenlarsen Jun 25, 2025

Uh oh!

dm-vodopyanov Jun 25, 2025

Uh oh!

steffenlarsen Jun 25, 2025

Uh oh!

Uh oh!

dm-vodopyanov Jun 24, 2025

Uh oh!

steffenlarsen Jun 25, 2025

Uh oh!

Uh oh!

Uh oh!

[SYCL] Implement loading SYCLBIN into kernel_bundle #18949

[SYCL] Implement loading SYCLBIN into kernel_bundle #18949

Uh oh!

Conversation

steffenlarsen commented Jun 12, 2025

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dm-vodopyanov left a comment

Choose a reason for hiding this comment

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!

Uh oh!

Uh oh!