[SYCL] Implement loading SYCLBIN into kernel_bundle (#18949)

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.

---------

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

Author: steffenlarsen

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

@@ -0,0 +1,82 @@
+//==---- syclbin_kernel_bundle.hpp - SYCLBIN-based kernel_bundle tooling ---==//
+//
+// 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>
+
+#include <fstream>
+#include <string>
+
+#if __has_include(<filesystem>)
+#include <filesystem>
+#endif
+
+#if __has_include(<span>)
+#include <span>
+#endif
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+template <bundle_state State, typename PropertyListT = empty_properties_t>
+std::enable_if_t<State != bundle_state::ext_oneapi_source, kernel_bundle<State>>
+get_kernel_bundle(const context &Ctxt, const std::vector<device> &Devs,
+                  const sycl::span<char> &Bytes, PropertyListT = {}) {
+  std::vector<device> UniqueDevices =
+      sycl::detail::removeDuplicateDevices(Devs);
+
+  sycl::detail::KernelBundleImplPtr Impl =
+      sycl::detail::get_kernel_bundle_impl(Ctxt, UniqueDevices, Bytes, State);
+  return sycl::detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
+}
+
+#if __cpp_lib_span
+template <bundle_state State, typename PropertyListT = empty_properties_t>
+std::enable_if_t<State != bundle_state::ext_oneapi_source, kernel_bundle<State>>
+get_kernel_bundle(const context &Ctxt, const std::vector<device> &Devs,
+                  const std::span<char> &Bytes, PropertyListT Props = {}) {
+  return experimental::get_kernel_bundle(
+      Ctxt, Devs, sycl::span<char>(Bytes.data(), Bytes.size()), Props);
+}
+#endif
+
+#if __cpp_lib_filesystem
+template <bundle_state State, typename PropertyListT = empty_properties_t>
+std::enable_if_t<State != bundle_state::ext_oneapi_source, kernel_bundle<State>>
+get_kernel_bundle(const context &Ctxt, const std::vector<device> &Devs,
+                  const std::filesystem::path &Filename,
+                  PropertyListT Props = {}) {
+  std::vector<char> RawSYCLBINData;
+  {
+    std::ifstream FileStream{Filename, std::ios::binary};
+    if (!FileStream.is_open())
+      throw std::ios_base::failure("Failed to open SYCLBIN file: " +
+                                   Filename.string());
+    RawSYCLBINData =
+        std::vector<char>{std::istreambuf_iterator<char>(FileStream),
+                          std::istreambuf_iterator<char>()};
+  }
+  return experimental::get_kernel_bundle<State>(
+      Ctxt, Devs, sycl::span<char>{RawSYCLBINData}, Props);
+}
+
+template <bundle_state State, typename PropertyListT = empty_properties_t>
+std::enable_if_t<State != bundle_state::ext_oneapi_source, kernel_bundle<State>>
+get_kernel_bundle(const context &Ctxt, const std::filesystem::path &Filename,
+                  PropertyListT Props = {}) {
+  return experimental::get_kernel_bundle<State>(Ctxt, Ctxt.get_devices(),
+                                                Filename, Props);
+}
+#endif
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/kernel_bundle.hpp

@@ -19,7 +19,8 @@
 #include <sycl/kernel.hpp>              // for kernel, kernel_bundle
 #include <sycl/kernel_bundle_enums.hpp> // for bundle_state
 #include <sycl/property_list.hpp>       // for property_list
-#include <ur_api.h>                     // for ur_native_handle_t
+#include <sycl/sycl_span.hpp>
+#include <ur_api.h>
 
 #include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>     // PropertyT
@@ -639,6 +640,10 @@ __SYCL_EXPORT detail::KernelBundleImplPtr
 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
                        bundle_state State);
 
+__SYCL_EXPORT detail::KernelBundleImplPtr
+get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
+                       const sycl::span<char> &Bytes, bundle_state State);
+
 __SYCL_EXPORT const std::vector<device>
 removeDuplicateDevices(const std::vector<device> &Devs);
 

sycl/include/sycl/sycl.hpp

@@ -105,6 +105,7 @@
 #include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
 #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/tangle_group.hpp>
 #include <sycl/ext/oneapi/experimental/work_group_memory.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>

sycl/source/CMakeLists.txt

@@ -301,6 +301,7 @@ set(SYCL_COMMON_SOURCES
     "detail/reduction.cpp"
     "detail/sampler_impl.cpp"
     "detail/stream_impl.cpp"
+    "detail/syclbin.cpp"
     "detail/scheduler/commands.cpp"
     "detail/scheduler/leaves_collection.cpp"
     "detail/scheduler/scheduler.cpp"

sycl/source/detail/base64.hpp

@@ -0,0 +1,121 @@
+//===--- Base64.h - Base64 Encoder/Decoder ----------------------*- C++ -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// Adjusted copy of llvm/include/llvm/Support/Base64.h.
+// TODO: Remove once we can consistently link the SYCL runtime library with
+// LLVMSupport.
+
+#pragma once
+
+#include <cstdint>
+#include <memory>
+#include <string>
+#include <vector>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+class Base64 {
+private:
+  // Decode a single character.
+  static inline int decode(char Ch) {
+    if (Ch >= 'A' && Ch <= 'Z') // 0..25
+      return Ch - 'A';
+    else if (Ch >= 'a' && Ch <= 'z') // 26..51
+      return Ch - 'a' + 26;
+    else if (Ch >= '0' && Ch <= '9') // 52..61
+      return Ch - '0' + 52;
+    else if (Ch == '+') // 62
+      return 62;
+    else if (Ch == '/') // 63
+      return 63;
+    return -1;
+  }
+
+  // Decode a quadruple of characters.
+  static inline void decode4(const char *Src, byte *Dst) {
+    int BadCh = -1;
+
+    for (auto I = 0; I < 4; ++I) {
+      char Ch = Src[I];
+      int Byte = decode(Ch);
+
+      if (Byte < 0) {
+        BadCh = Ch;
+        break;
+      }
+      Dst[I] = (byte)Byte;
+    }
+    if (BadCh != -1)
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "Invalid char in base 64 encoding.");
+  }
+
+public:
+  using byte = uint8_t;
+
+  // Get the size of the encoded byte sequence of given size.
+  static size_t getDecodedSize(size_t SrcSize) { return (SrcSize * 3 + 3) / 4; }
+
+  // Decode a sequence of given size into a pre-allocated memory.
+  // Returns the number of bytes in the decoded result or 0 in case of error.
+  static size_t decode(const char *Src, byte *Dst, size_t SrcSize) {
+    size_t SrcOff = 0;
+    size_t DstOff = 0;
+
+    // decode full quads
+    for (size_t Qch = 0; Qch < SrcSize / 4; ++Qch, SrcOff += 4, DstOff += 3) {
+      byte Ch[4] = {0, 0, 0, 0};
+      decode4(Src + SrcOff, Ch);
+
+      // each quad of chars produces three bytes of output
+      Dst[DstOff + 0] = Ch[0] | (Ch[1] << 6);
+      Dst[DstOff + 1] = (Ch[1] >> 2) | (Ch[2] << 4);
+      Dst[DstOff + 2] = (Ch[2] >> 4) | (Ch[3] << 2);
+    }
+    auto RemChars = SrcSize - SrcOff;
+
+    if (RemChars == 0)
+      return DstOff;
+    // decode the remainder; variants:
+    // 2 chars remain - produces single byte
+    // 3 chars remain - produces two bytes
+
+    if (RemChars != 2 && RemChars != 3)
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "Invalid encoded sequence length.");
+
+    int Ch0 = decode(Src[SrcOff++]);
+    int Ch1 = decode(Src[SrcOff++]);
+    int Ch2 = RemChars == 3 ? decode(Src[SrcOff]) : 0;
+
+    if (Ch0 < 0 || Ch1 < 0 || Ch2 < 0)
+      throw sycl::exception(
+          make_error_code(errc::invalid),
+          "Invalid characters in the encoded sequence remainder.");
+    Dst[DstOff++] = Ch0 | (Ch1 << 6);
+
+    if (RemChars == 3)
+      Dst[DstOff++] = (Ch1 >> 2) | (Ch2 << 4);
+    return DstOff;
+  }
+
+  // Allocate minimum required amount of memory and decode a sequence of given
+  // size into it.
+  // Returns the decoded result. The size can be obtained via getDecodedSize.
+  static std::unique_ptr<byte[]> decode(const char *Src, size_t SrcSize) {
+    size_t DstSize = getDecodedSize(SrcSize);
+    std::unique_ptr<byte[]> Dst(new byte[DstSize]);
+    decode(Src, Dst.get(), SrcSize);
+    return Dst;
+  }
+};
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

sycl/source/detail/compiler.hpp

@@ -57,6 +57,8 @@
 #define __SYCL_PROPERTY_SET_SYCL_MISC_PROP "SYCL/misc properties"
 /// PropertySetRegistry::SYCL_ASSERT_USED defined in PropertySetIO.h
 #define __SYCL_PROPERTY_SET_SYCL_ASSERT_USED "SYCL/assert used"
+/// PropertySetRegistry::SYCL_KERNEL_NAMES defined in PropertySetIO.h
+#define __SYCL_PROPERTY_SET_SYCL_KERNEL_NAMES "SYCL/kernel names"
 /// PropertySetRegistry::SYCL_EXPORTED_SYMBOLS defined in PropertySetIO.h
 #define __SYCL_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS "SYCL/exported symbols"
 /// PropertySetRegistry::SYCL_IMPORTED_SYMBOLS defined in PropertySetIO.h

sycl/source/detail/device_binary_image.cpp

@@ -162,7 +162,7 @@ RTDeviceBinaryImage::getProperty(const char *PropName) const {
   return *It;
 }
 
-void RTDeviceBinaryImage::init(sycl_device_binary Bin) {
+RTDeviceBinaryImage::RTDeviceBinaryImage(sycl_device_binary Bin) {
   ImageId = ImageCounter++;
 
   // If there was no binary, we let the owner handle initialization as they see
@@ -199,6 +199,7 @@ void RTDeviceBinaryImage::init(sycl_device_binary Bin) {
     ProgramMetadataUR.push_back(
         ur::mapDeviceBinaryPropertyToProgramMetadata(Prop));
   }
+  KernelNames.init(Bin, __SYCL_PROPERTY_SET_SYCL_KERNEL_NAMES);
   ExportedSymbols.init(Bin, __SYCL_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS);
   ImportedSymbols.init(Bin, __SYCL_PROPERTY_SET_SYCL_IMPORTED_SYMBOLS);
   DeviceGlobals.init(Bin, __SYCL_PROPERTY_SET_SYCL_DEVICE_GLOBALS);
@@ -211,7 +212,8 @@ void RTDeviceBinaryImage::init(sycl_device_binary Bin) {
 
 std::atomic<uintptr_t> RTDeviceBinaryImage::ImageCounter = 1;
 
-DynRTDeviceBinaryImage::DynRTDeviceBinaryImage() : RTDeviceBinaryImage() {
+DynRTDeviceBinaryImage::DynRTDeviceBinaryImage()
+    : RTDeviceBinaryImage(nullptr) {
   Bin = new sycl_device_binary_struct();
   Bin->Version = SYCL_DEVICE_BINARY_VERSION;
   Bin->Kind = SYCL_DEVICE_BINARY_OFFLOAD_KIND_SYCL;
@@ -227,12 +229,11 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage() : RTDeviceBinaryImage() {
   Bin->DeviceTargetSpec = __SYCL_DEVICE_BINARY_TARGET_UNKNOWN;
 }
 
-DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
-    std::unique_ptr<char[], std::function<void(void *)>> &&DataPtr,
-    size_t DataSize)
-    : DynRTDeviceBinaryImage() {
-  Data = std::move(DataPtr);
-  Bin->BinaryStart = reinterpret_cast<unsigned char *>(Data.get());
+std::unique_ptr<sycl_device_binary_struct> CreateDefaultDynBinary(
+    const std::unique_ptr<char[], std::function<void(void *)>> &DataPtr,
+    size_t DataSize) {
+  auto Bin = std::make_unique<sycl_device_binary_struct>();
+  Bin->BinaryStart = reinterpret_cast<unsigned char *>(DataPtr.get());
   Bin->BinaryEnd = Bin->BinaryStart + DataSize;
   Bin->Format = ur::getBinaryImageFormat(Bin->BinaryStart, DataSize);
   switch (Bin->Format) {
@@ -242,9 +243,15 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
   default:
     Bin->DeviceTargetSpec = __SYCL_DEVICE_BINARY_TARGET_UNKNOWN;
   }
-  init(Bin);
+  return Bin;
 }
 
+DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
+    std::unique_ptr<char[], std::function<void(void *)>> &&DataPtr,
+    size_t DataSize)
+    : RTDeviceBinaryImage(CreateDefaultDynBinary(DataPtr, DataSize).release()),
+      Data{std::move(DataPtr)} {}
+
 DynRTDeviceBinaryImage::~DynRTDeviceBinaryImage() {
   delete Bin;
   Bin = nullptr;
@@ -479,8 +486,6 @@ static void copyProperty(sycl_device_binary_property &NextFreeProperty,
 DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
     const std::vector<const RTDeviceBinaryImage *> &Imgs)
     : DynRTDeviceBinaryImage() {
-  init(nullptr);
-
   // Naive merges.
   auto MergedSpecConstants =
       naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) {
@@ -510,6 +515,10 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
       naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) {
         return Img.getImplicitLocalArg();
       });
+  auto MergedKernelNames =
+      naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) {
+        return Img.getKernelNames();
+      });
   auto MergedExportedSymbols =
       naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) {
         return Img.getExportedSymbols();
@@ -519,12 +528,13 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
         return Img.getRegisteredKernels();
       });
 
-  std::array<const std::vector<sycl_device_binary_property> *, 10> MergedVecs{
+  std::array<const std::vector<sycl_device_binary_property> *, 11> MergedVecs{
       &MergedSpecConstants,      &MergedSpecConstantsDefaultValues,
       &MergedKernelParamOptInfo, &MergedAssertUsed,
       &MergedDeviceGlobals,      &MergedHostPipes,
       &MergedVirtualFunctions,   &MergedImplicitLocalArg,
-      &MergedExportedSymbols,    &MergedRegisteredKernels};
+      &MergedKernelNames,        &MergedExportedSymbols,
+      &MergedRegisteredKernels};
 
   // Exclusive merges.
   auto MergedDeviceLibReqMask =
@@ -648,6 +658,7 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
   CopyPropertiesVec(MergedHostPipes, HostPipes);
   CopyPropertiesVec(MergedVirtualFunctions, VirtualFunctions);
   CopyPropertiesVec(MergedImplicitLocalArg, ImplicitLocalArg);
+  CopyPropertiesVec(MergedKernelNames, KernelNames);
   CopyPropertiesVec(MergedExportedSymbols, ExportedSymbols);
   CopyPropertiesVec(MergedRegisteredKernels, RegisteredKernels);
 
@@ -675,18 +686,11 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
 #ifdef SYCL_RT_ZSTD_AVAILABLE
 CompressedRTDeviceBinaryImage::CompressedRTDeviceBinaryImage(
     sycl_device_binary CompressedBin)
-    : RTDeviceBinaryImage() {
-
-  // 'CompressedBin' is part of the executable image loaded into memory
-  // which can't be modified easily. So, we need to make a copy of it.
-  Bin = new sycl_device_binary_struct(*CompressedBin);
-
+    : RTDeviceBinaryImage(new sycl_device_binary_struct(*CompressedBin)) {
   // Get the decompressed size of the binary image.
   m_ImageSize = ZSTDCompressor::GetDecompressedSize(
       reinterpret_cast<const char *>(Bin->BinaryStart),
       static_cast<size_t>(Bin->BinaryEnd - Bin->BinaryStart));
-
-  init(Bin);
 }
 
 void CompressedRTDeviceBinaryImage::Decompress() {