[llvm] [libsycl] Add liboffload kernel creation (PR #188794)
Kseniya Tikhomirova via llvm-commits
llvm-commits at lists.llvm.org
Tue Apr 14 02:20:42 PDT 2026
https://github.com/KseniyaTikhomirova updated https://github.com/llvm/llvm-project/pull/188794
>From 913cb645cd56e4bd9d56b28d76809d099091a643 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Thu, 19 Mar 2026 05:17:36 -0700
Subject: [PATCH 1/9] [libsycl] Add device image registration
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/docs/index.rst | 5 +
libsycl/src/CMakeLists.txt | 1 +
.../src/detail/device_binary_structures.hpp | 129 +++++++++++++++
libsycl/src/detail/device_image_wrapper.hpp | 51 ++++++
libsycl/src/detail/device_impl.hpp | 2 +-
libsycl/src/detail/global_objects.cpp | 3 -
libsycl/src/detail/global_objects.hpp | 13 --
libsycl/src/detail/kernel_id.hpp | 77 +++++++++
libsycl/src/detail/program_manager.cpp | 154 ++++++++++++++++++
libsycl/src/detail/program_manager.hpp | 95 +++++++++++
libsycl/src/usm_functions.cpp | 2 +-
11 files changed, 514 insertions(+), 18 deletions(-)
create mode 100644 libsycl/src/detail/device_binary_structures.hpp
create mode 100644 libsycl/src/detail/device_image_wrapper.hpp
create mode 100644 libsycl/src/detail/kernel_id.hpp
create mode 100644 libsycl/src/detail/program_manager.cpp
create mode 100644 libsycl/src/detail/program_manager.hpp
diff --git a/libsycl/docs/index.rst b/libsycl/docs/index.rst
index 0ec3a4d825075..03f7fb7c0876e 100644
--- a/libsycl/docs/index.rst
+++ b/libsycl/docs/index.rst
@@ -113,3 +113,8 @@ TODO for added SYCL classes
* add aligned functions (blocked by liboffload support)
* forward templated funcs to alignment methods (rewrite current impl)
* handle sub devices once they are implemented (blocked by liboffload support)
+
+
+* general opens:
+
+ * define a way to report errors from object dtors.
\ No newline at end of file
diff --git a/libsycl/src/CMakeLists.txt b/libsycl/src/CMakeLists.txt
index 67ba7d28968de..bc39f9162f992 100644
--- a/libsycl/src/CMakeLists.txt
+++ b/libsycl/src/CMakeLists.txt
@@ -93,6 +93,7 @@ set(LIBSYCL_SOURCES
"detail/device_impl.cpp"
"detail/global_objects.cpp"
"detail/platform_impl.cpp"
+ "detail/program_manager.cpp"
"detail/queue_impl.cpp"
"detail/offload/offload_utils.cpp"
"detail/offload/offload_topology.cpp"
diff --git a/libsycl/src/detail/device_binary_structures.hpp b/libsycl/src/detail/device_binary_structures.hpp
new file mode 100644
index 0000000000000..d2abac6854e40
--- /dev/null
+++ b/libsycl/src/detail/device_binary_structures.hpp
@@ -0,0 +1,129 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_DEVICE_BINARY_STRUCTURES
+#define _LIBSYCL_DEVICE_BINARY_STRUCTURES
+
+#include <cstdint>
+
+/// Target identification strings.
+///
+/// A device type represented by a particular target
+/// triple requires specific binary images. We need
+/// to map the image type onto the device target triple.
+
+/// Unknown target.
+#define _LIBSYCL_DEVICE_BINARY_TARGET_TRIPLE_UNKNOWN "unknown-unknown-unknown"
+/// SPIR-V with 64-bit pointers.
+#define _LIBSYCL_DEVICE_BINARY_TARGET_SPIRV64 "spirv64-unknown-unknown"
+
+// This is a replica of the EntryTy data structure in
+// llvm/include/llvm/Frontend/Offloading/Utility.h.
+struct EntryTy {
+ /// Reserved bytes used to detect an older version of the struct, always zero.
+ uint64_t Reserved = 0x0;
+ /// The current version of the struct for runtime forward compatibility.
+ uint16_t Version = 0x1;
+ /// The expected consumer of this entry, e.g. SYCL or OpenMP.
+ /// llvm::object::OffloadKind
+ uint16_t Kind;
+ /// Flags associated with the global.
+ uint32_t Flags;
+ /// The address of the global to be registered by the runtime.
+ void *Address;
+ /// The name of the symbol in the device image.
+ char *SymbolName;
+ /// The number of bytes the symbol takes.
+ uint64_t Size;
+ /// Extra generic data used to register this entry.
+ uint64_t Data;
+ /// An extra pointer, usually null.
+ void *AuxAddr;
+};
+
+// TODO: would be nice to include it from llvm/Object. It doesn't work now since
+// I have to link LLVMObject to do that (linker error otherwise).
+// Copy of llvm::object::OffloadKind.
+/// The producer of the associated offloading image.
+enum OffloadKind : uint16_t {
+ OFK_None = 0,
+ OFK_OpenMP = (1 << 0),
+ OFK_Cuda = (1 << 1),
+ OFK_HIP = (1 << 2),
+ OFK_SYCL = (1 << 3),
+ OFK_LAST = (1 << 4),
+};
+
+// Copy of llvm::object::ImageKind.
+/// The type of contents the offloading image contains.
+enum ImageKind : uint16_t {
+ IMG_None = 0,
+ IMG_Object,
+ IMG_Bitcode,
+ IMG_Cubin,
+ IMG_Fatbinary,
+ IMG_PTX,
+ IMG_SPIRV,
+ IMG_LAST,
+};
+
+/// Device binary descriptor version supported by this library.
+static const uint16_t _LIBSYCL_SUPPORTED_DEVICE_BINARY_VERSION = 3;
+
+/// This struct is a record of the device binary information.
+/// It must match the __tgt_device_image structure generated by the
+/// clang-offload-wrapper tool when their `Version` field match.
+struct __sycl_tgt_device_image {
+ uint16_t Version;
+ /// The type of offload model the binary employs. See `OffloadKind`. Only
+ /// OFK_SYCL is supported by libsycl.
+ uint8_t OffloadKind;
+ /// Format of the binary data, see `ImageKind`.
+ uint8_t ImageFormat;
+ /// A null-terminated string representation of the device's target
+ /// architecture. Must hold one of _LIBSYCL_DEVICE_BINARY_TARGET_* values.
+ const char *TripleString;
+ /// A null-terminated string; target- and compiler-specific options
+ /// which are suggested to use to "compile" program at runtime.
+ const char *CompileOptions;
+ /// A null-terminated string; target- and compiler-specific options
+ /// which are suggested to use to "link" program at runtime.
+ const char *LinkOptions;
+ /// Pointer to the target code start.
+ const unsigned char *ImageStart;
+ /// Pointer to the target code end.
+ const unsigned char *ImageEnd;
+ /// The offload entry table
+ EntryTy *EntriesBegin;
+ EntryTy *EntriesEnd;
+ // TODO: properties are not supported now.
+ /// Array of preperty sets.
+ void *PropertiesBegin;
+ void *PropertiesEnd;
+};
+
+/// Version of offload binaries descriptor `__sycl_tgt_bin_desc` supported by
+/// libsycl.
+static constexpr uint16_t _LIBSYCL_SUPPORTED_OFFLOAD_BINARY_VERSION = 1;
+
+/// This struct is a record of all the device code that may be offloaded.
+/// It must match the `__tgt_bin_desc` structure generated by
+/// the clang-offload-wrapper tool when their `Version` field match.
+struct __sycl_tgt_bin_desc {
+ /// Version of the structure.
+ uint16_t Version;
+ /// Number of device binaries in this descriptor.
+ uint16_t NumDeviceBinaries;
+ /// Device binaries data.
+ __sycl_tgt_device_image *DeviceImages;
+ /// The offload entry table (not used, for compatibility with OpenMP).
+ EntryTy *HostEntriesBegin;
+ EntryTy *HostEntriesEnd;
+};
+
+#endif // _LIBSYCL_DEVICE_BINARY_STRUCTURES
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
new file mode 100644
index 0000000000000..534a4be4fcf1e
--- /dev/null
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -0,0 +1,51 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_DEVICE_IMAGE_WRAPPER
+#define _LIBSYCL_DEVICE_IMAGE_WRAPPER
+
+#include <sycl/__impl/detail/config.hpp>
+
+#include <detail/device_binary_structures.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+namespace detail {
+
+/// A wrapper of __sycl_tgt_device_image structure.
+class DeviceImageWrapper {
+public:
+public:
+ DeviceImageWrapper(const __sycl_tgt_device_image &Bin) : MBin(&Bin) {}
+ // Explicitly delete copy constructor/operator= to avoid unintentional copies.
+ DeviceImageWrapper(const DeviceImageWrapper &) = delete;
+ DeviceImageWrapper &operator=(const DeviceImageWrapper &) = delete;
+
+ DeviceImageWrapper(DeviceImageWrapper &&) = default;
+ DeviceImageWrapper &operator=(DeviceImageWrapper &&) = default;
+
+ ~DeviceImageWrapper() {}
+
+ /// \return reference to corresponsing raw __sycl_tgt_device_image object.
+ const __sycl_tgt_device_image &getRawData() const { return *get(); }
+
+ /// \return size of corresponding device image data in bytes.
+ size_t getSize() const {
+ return static_cast<size_t>(MBin->ImageEnd - MBin->ImageStart);
+ }
+
+protected:
+ const __sycl_tgt_device_image *get() const { return MBin; }
+
+ __sycl_tgt_device_image const *MBin{};
+};
+
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL_DEVICE_IMAGE_WRAPPER
diff --git a/libsycl/src/detail/device_impl.hpp b/libsycl/src/detail/device_impl.hpp
index c83b767aad02f..3c96166be2e8b 100644
--- a/libsycl/src/detail/device_impl.hpp
+++ b/libsycl/src/detail/device_impl.hpp
@@ -115,7 +115,7 @@ class DeviceImpl {
static_assert(false && "Info descriptor is not properly supported");
}
- ol_device_handle_t getOLHandle() { return MOffloadDevice; }
+ ol_device_handle_t getHandle() { return MOffloadDevice; }
private:
ol_device_handle_t MOffloadDevice = {};
diff --git a/libsycl/src/detail/global_objects.cpp b/libsycl/src/detail/global_objects.cpp
index d80be710268f8..35e32985e7cbb 100644
--- a/libsycl/src/detail/global_objects.cpp
+++ b/libsycl/src/detail/global_objects.cpp
@@ -53,6 +53,3 @@ std::vector<PlatformImplUPtr> &getPlatformCache() {
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
-
-extern "C" void __sycl_register_lib(void *) {}
-extern "C" void __sycl_unregister_lib(void *) {}
diff --git a/libsycl/src/detail/global_objects.hpp b/libsycl/src/detail/global_objects.hpp
index 008cb01f4f355..4535a254c6609 100644
--- a/libsycl/src/detail/global_objects.hpp
+++ b/libsycl/src/detail/global_objects.hpp
@@ -16,19 +16,6 @@
#include <mutex>
#include <vector>
-// +++ Entry points referenced by the offload wrapper object {
-
-/// Executed as a part of current module's (.exe, .dll) static initialization.
-/// Registers device executable images with the runtime.
-extern "C" _LIBSYCL_EXPORT void __sycl_register_lib(void *);
-
-/// Executed as a part of current module's (.exe, .dll) static
-/// de-initialization.
-/// Unregisters device executable images with the runtime.
-extern "C" _LIBSYCL_EXPORT void __sycl_unregister_lib(void *);
-
-// +++ }
-
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
diff --git a/libsycl/src/detail/kernel_id.hpp b/libsycl/src/detail/kernel_id.hpp
new file mode 100644
index 0000000000000..a1508042ed960
--- /dev/null
+++ b/libsycl/src/detail/kernel_id.hpp
@@ -0,0 +1,77 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_KERNEL_ID
+#define _LIBSYCL_KERNEL_ID
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/detail/obj_utils.hpp>
+
+#include <memory>
+#include <string>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+/// The class is impl counterpart for sycl::kernel_id which represent a kernel
+/// identificator.
+class KernelIdImpl {
+public:
+ KernelIdImpl(std::string_view Name) : MName(std::move(Name)) {}
+ KernelIdImpl() {}
+ /// \return a null-terminated string representing the name of kernel this id
+ /// stands for.
+ const char *get_name() { return MName.data(); }
+
+private:
+ std::string MName;
+};
+} // namespace detail
+
+// TODO: It is not exported now, but is a part of SYCL spec.
+/// Kernel identifier.
+class kernel_id {
+public:
+ kernel_id() = delete;
+
+ kernel_id(const kernel_id &rhs) = default;
+
+ kernel_id(kernel_id &&rhs) = default;
+
+ kernel_id &operator=(const kernel_id &rhs) = default;
+
+ kernel_id &operator=(kernel_id &&rhs) = default;
+
+ friend bool operator==(const kernel_id &lhs, const kernel_id &rhs) {
+ return lhs.impl == rhs.impl;
+ }
+
+ friend bool operator!=(const kernel_id &lhs, const kernel_id &rhs) {
+ return !(lhs == rhs);
+ }
+
+ /// \returns a null-terminated string which contains the kernel name.
+ const char *get_name() const noexcept;
+
+private:
+ kernel_id(const char *Name);
+
+ kernel_id(const std::shared_ptr<detail::KernelIdImpl> &Impl)
+ : impl(std::move(Impl)) {}
+
+ std::shared_ptr<detail::KernelIdImpl> impl;
+ friend sycl::detail::ImplUtils;
+};
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+template <>
+struct std::hash<sycl::kernel_id>
+ : public sycl::detail::HashBase<sycl::kernel_id> {};
+
+#endif // _LIBSYCL_KERNEL_ID
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
new file mode 100644
index 0000000000000..ea8673e0f39e9
--- /dev/null
+++ b/libsycl/src/detail/program_manager.cpp
@@ -0,0 +1,154 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <sycl/__impl/exception.hpp>
+
+#include <detail/device_impl.hpp>
+#include <detail/offload/offload_utils.hpp>
+#include <detail/program_manager.hpp>
+
+#include <cstring>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+namespace detail {
+
+static inline bool checkFatBinVersion(const __sycl_tgt_bin_desc &FatbinDesc) {
+ return FatbinDesc.Version == _LIBSYCL_SUPPORTED_OFFLOAD_BINARY_VERSION;
+}
+
+static inline bool
+checkDeviceImageValidness(const __sycl_tgt_device_image &DeviceImage) {
+ return (DeviceImage.Version == _LIBSYCL_SUPPORTED_DEVICE_BINARY_VERSION) &&
+ (DeviceImage.OffloadKind == OFK_SYCL) &&
+ (DeviceImage.ImageFormat == IMG_SPIRV);
+}
+
+void ProgramManager::addImages(__sycl_tgt_bin_desc *FatbinDesc) {
+ assert(FatbinDesc && "Device images descriptor can't be nullptr");
+
+ if (!checkFatBinVersion(*FatbinDesc))
+ throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+ "Incompatible version of device images descriptor.");
+ if (!FatbinDesc->NumDeviceBinaries)
+ return;
+
+ std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
+ for (int I = 0; I < FatbinDesc->NumDeviceBinaries; I++) {
+ const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
+ if (!checkDeviceImageValidness(RawDeviceImage))
+ throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+ "Incompatible device image.");
+
+ const EntryTy *EntriesB = RawDeviceImage.EntriesBegin;
+ const EntryTy *EntriesE = RawDeviceImage.EntriesEnd;
+ // Ignore "empty" device image.
+ if (EntriesB == EntriesE)
+ continue;
+
+ std::unique_ptr<DeviceImageWrapper> NewImageWrapper =
+ std::make_unique<DeviceImageWrapper>(RawDeviceImage);
+
+ for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt++) {
+ auto Name = EntriesIt->SymbolName;
+ auto KernelIDIt = MKernelNameToID.find(Name);
+ if (KernelIDIt == MKernelNameToID.end()) {
+ sycl::kernel_id KernelID =
+ detail::createSyclObjFromImpl<sycl::kernel_id>(
+ std::make_shared<detail::KernelIdImpl>(Name));
+ KernelIDIt = MKernelNameToID.insert(
+ MKernelNameToID.end(),
+ std::make_pair(std::string_view(Name), KernelID));
+ }
+
+ MKernelIDToDevImageJIT.insert(
+ std::make_pair(KernelIDIt->second, NewImageWrapper.get()));
+ }
+
+ MDeviceImageWrappers.insert(
+ std::make_pair(&RawDeviceImage, std::move(NewImageWrapper)));
+ }
+}
+
+void ProgramManager::removeImages(__sycl_tgt_bin_desc *FatbinDesc) {
+ assert(FatbinDesc && "Device images descriptor can't be nullptr");
+
+ if (!checkFatBinVersion(*FatbinDesc))
+ throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+ "Incompatible version of device images descriptor.");
+ if (FatbinDesc->NumDeviceBinaries == 0)
+ return;
+
+ std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
+ for (int I = 0; I < FatbinDesc->NumDeviceBinaries; I++) {
+ const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
+
+ auto DevImageIt = MDeviceImageWrappers.find(&RawDeviceImage);
+ if (DevImageIt == MDeviceImageWrappers.end())
+ continue;
+
+ const EntryTy *EntriesB = RawDeviceImage.EntriesBegin;
+ const EntryTy *EntriesE = RawDeviceImage.EntriesEnd;
+ // Ignore "empty" device image
+ if (EntriesB == EntriesE)
+ continue;
+
+ for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt++) {
+ if (auto KernelIDIt = MKernelNameToID.find(EntriesIt->SymbolName);
+ KernelIDIt != MKernelNameToID.end()) {
+ MKernelIDToDevImageJIT.erase(KernelIDIt->second);
+ MKernelNameToID.erase(KernelIDIt);
+ }
+ }
+
+ MDeviceImageWrappers.erase(DevImageIt);
+ }
+}
+
+static bool isImageTargetCompatible(const DeviceImageWrapper &Image,
+ const DeviceImpl &Device) {
+ sycl::backend BE = Device.getBackend();
+ const char *Target = Image.getRawData().TripleString;
+
+ return (strcmp(Target, _LIBSYCL_DEVICE_BINARY_TARGET_SPIRV64) == 0) &&
+ (BE == sycl::backend::level_zero);
+}
+
+DeviceImageWrapper *ProgramManager::getDeviceImage(const char *KernelName,
+ kernel_id KernelID,
+ DeviceImpl &Device) {
+ auto [Begin, End] = MKernelIDToDevImageJIT.equal_range(KernelID);
+ if (Begin != End) {
+ ol_result_t Result{};
+ bool IsValid{};
+ // TODO: with AOT (not implemented yet), we need to analize and check
+ // olIsValidBinary for AOT binaries first.
+ for (auto It = Begin; It != End; ++It) {
+ if (isImageTargetCompatible(*It->second, Device)) {
+ callAndThrow(olIsValidBinary, Device.getHandle(),
+ It->second->getRawData().ImageStart, It->second->getSize(),
+ &IsValid);
+ if (IsValid)
+ return It->second;
+ }
+ }
+ }
+
+ throw exception(make_error_code(errc::runtime),
+ "No kernel named " + std::string(KernelName) + " was found");
+}
+
+} // namespace detail
+_LIBSYCL_END_NAMESPACE_SYCL
+
+extern "C" void __sycl_register_lib(__sycl_tgt_bin_desc *FatbinDesc) {
+ sycl::detail::ProgramManager::getInstance().addImages(FatbinDesc);
+}
+
+extern "C" void __sycl_unregister_lib(__sycl_tgt_bin_desc *FatbinDesc) {
+ sycl::detail::ProgramManager::getInstance().removeImages(FatbinDesc);
+}
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
new file mode 100644
index 0000000000000..3b1a2eb07d5ff
--- /dev/null
+++ b/libsycl/src/detail/program_manager.hpp
@@ -0,0 +1,95 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_PROGRAM_MANAGER
+#define _LIBSYCL_PROGRAM_MANAGER
+
+#include <sycl/__impl/detail/config.hpp>
+
+#include <detail/device_binary_structures.hpp>
+#include <detail/device_image_wrapper.hpp>
+#include <detail/kernel_id.hpp>
+
+#include <mutex>
+#include <unordered_map>
+
+// +++ Entry points referenced by the offload wrapper object {
+
+/// Executed as a part of current module's (.exe, .dll) static initialization.
+/// Registers device executable images with the runtime.
+extern "C" _LIBSYCL_EXPORT void __sycl_register_lib(__sycl_tgt_bin_desc *desc);
+
+/// Executed as a part of current module's (.exe, .dll) static
+/// de-initialization.
+/// Unregisters device executable images with the runtime.
+extern "C" _LIBSYCL_EXPORT void
+__sycl_unregister_lib(__sycl_tgt_bin_desc *desc);
+
+// +++ }
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+class DeviceImpl;
+
+/// A class to manage programs and kernels.
+class ProgramManager {
+
+public:
+ static ProgramManager &getInstance() {
+ static ProgramManager PM{};
+ return PM;
+ }
+
+ /// Parses raw device images data and prepare internal structures for
+ /// effective kernel/program creation.
+ /// \param FatbinDesc a record of all the device code that may be offloaded
+ /// generated by compiler and offloading tools.
+ /// \throw sycl::exception with sycl::errc::runtime if device image descriptor
+ /// has incompatible version or if device image has incompatible
+ /// version/target/kind.
+ void addImages(__sycl_tgt_bin_desc *FatbinDesc);
+
+ /// Removes all entries of the data in FatbinDesc in internal structures.
+ /// \param FatbinDesc a record of all the device code that may be offloaded
+ /// generated by compiler and offloading tools. Must match the pointer and
+ /// data passed to addImages.
+ void removeImages(__sycl_tgt_bin_desc *FatbinDesc);
+
+private:
+ ProgramManager() = default;
+ ~ProgramManager() = default;
+ ProgramManager(ProgramManager const &) = delete;
+ ProgramManager &operator=(ProgramManager const &) = delete;
+
+ /// Searches for device image that contains requested kernel and is compatible
+ /// with requested device.
+ /// \param KernelName a null-terminated string representing the name of kernel
+ /// to obtain device image for.
+ /// \param KernelID a kernel id matching KernelName.
+ /// \param DeviceImpl a device that device image must be compatible with.
+ /// \throw sycl::exception with sycl::errc::runtime if device image validness
+ /// check failed in liboffload or if no compatible image was found.
+ DeviceImageWrapper *getDeviceImage(const char *KernelName, kernel_id KernelID,
+ DeviceImpl &Device);
+
+ // Filled by addImages(...).
+ std::unordered_map<std::string_view, kernel_id> MKernelNameToID;
+ std::unordered_map<kernel_id, DeviceImageWrapper *> MKernelIDToDevImageJIT;
+ // Controls lifetime of device image ptr and wrapper.
+ std::unordered_map<const __sycl_tgt_device_image *,
+ std::unique_ptr<DeviceImageWrapper>>
+ MDeviceImageWrappers;
+ std::mutex MImageCollectionMutex;
+};
+
+} // namespace detail
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL_PROGRAM_MANAGER
diff --git a/libsycl/src/usm_functions.cpp b/libsycl/src/usm_functions.cpp
index 24a099ea4cf2a..c94015b97d772 100644
--- a/libsycl/src/usm_functions.cpp
+++ b/libsycl/src/usm_functions.cpp
@@ -103,7 +103,7 @@ void *malloc(std::size_t numBytes, const device &syclDevice,
void *Ptr{};
auto Result = detail::callNoCheck(
- olMemAlloc, detail::getSyclObjImpl(syclDevice)->getOLHandle(),
+ olMemAlloc, detail::getSyclObjImpl(syclDevice)->getHandle(),
detail::getOlAllocType(kind), numBytes, &Ptr);
return detail::isFailed(Result) ? nullptr : Ptr;
}
>From 87725e76b89cb8f56e6ff5a3b9e81b5ff10b55f7 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Tue, 24 Mar 2026 10:42:48 -0700
Subject: [PATCH 2/9] fix comments
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/CMakeLists.txt | 3 +
.../src/detail/device_binary_structures.hpp | 92 ++++++-------------
libsycl/src/detail/device_image_wrapper.hpp | 11 ++-
libsycl/src/detail/kernel_id.hpp | 14 +--
libsycl/src/detail/program_manager.cpp | 52 +++++------
libsycl/src/detail/program_manager.hpp | 38 ++++----
libsycl/src/detail/queue_impl.cpp | 3 +-
7 files changed, 91 insertions(+), 122 deletions(-)
diff --git a/libsycl/src/CMakeLists.txt b/libsycl/src/CMakeLists.txt
index bc39f9162f992..4501005e433e3 100644
--- a/libsycl/src/CMakeLists.txt
+++ b/libsycl/src/CMakeLists.txt
@@ -17,6 +17,7 @@ function(add_sycl_rt_library LIB_TARGET_NAME LIB_OBJ_NAME LIB_OUTPUT_NAME)
add_dependencies(${LIB_OBJ_NAME}
sycl-headers
LLVMOffload
+ LLVMObject
)
target_include_directories(${LIB_OBJ_NAME}
@@ -24,6 +25,7 @@ function(add_sycl_rt_library LIB_TARGET_NAME LIB_OBJ_NAME LIB_OUTPUT_NAME)
${CMAKE_CURRENT_SOURCE_DIR}
${LIBSYCL_BUILD_INCLUDE_DIR}
$<TARGET_PROPERTY:LLVMOffload,INTERFACE_INCLUDE_DIRECTORIES>
+ ${LLVM_MAIN_INCLUDE_DIR}
)
set_target_properties(${LIB_TARGET_NAME}
@@ -68,6 +70,7 @@ function(add_sycl_rt_library LIB_TARGET_NAME LIB_OBJ_NAME LIB_OUTPUT_NAME)
${CMAKE_DL_LIBS}
${CMAKE_THREAD_LIBS_INIT}
LLVMOffload
+ LLVMObject
)
set_target_properties(${LIB_TARGET_NAME}
diff --git a/libsycl/src/detail/device_binary_structures.hpp b/libsycl/src/detail/device_binary_structures.hpp
index d2abac6854e40..be403da9fcbb4 100644
--- a/libsycl/src/detail/device_binary_structures.hpp
+++ b/libsycl/src/detail/device_binary_structures.hpp
@@ -9,8 +9,17 @@
#ifndef _LIBSYCL_DEVICE_BINARY_STRUCTURES
#define _LIBSYCL_DEVICE_BINARY_STRUCTURES
+#include <sycl/__impl/detail/config.hpp>
+
+#include <llvm/Frontend/Offloading/Utility.h>
+#include <llvm/Object/OffloadBinary.h>
+
#include <cstdint>
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
/// Target identification strings.
///
/// A device type represented by a particular target
@@ -18,66 +27,16 @@
/// to map the image type onto the device target triple.
/// Unknown target.
-#define _LIBSYCL_DEVICE_BINARY_TARGET_TRIPLE_UNKNOWN "unknown-unknown-unknown"
+static constexpr char DeviceBinaryTripleUnknown[] = "unknown-unknown-unknown";
/// SPIR-V with 64-bit pointers.
-#define _LIBSYCL_DEVICE_BINARY_TARGET_SPIRV64 "spirv64-unknown-unknown"
-
-// This is a replica of the EntryTy data structure in
-// llvm/include/llvm/Frontend/Offloading/Utility.h.
-struct EntryTy {
- /// Reserved bytes used to detect an older version of the struct, always zero.
- uint64_t Reserved = 0x0;
- /// The current version of the struct for runtime forward compatibility.
- uint16_t Version = 0x1;
- /// The expected consumer of this entry, e.g. SYCL or OpenMP.
- /// llvm::object::OffloadKind
- uint16_t Kind;
- /// Flags associated with the global.
- uint32_t Flags;
- /// The address of the global to be registered by the runtime.
- void *Address;
- /// The name of the symbol in the device image.
- char *SymbolName;
- /// The number of bytes the symbol takes.
- uint64_t Size;
- /// Extra generic data used to register this entry.
- uint64_t Data;
- /// An extra pointer, usually null.
- void *AuxAddr;
-};
-
-// TODO: would be nice to include it from llvm/Object. It doesn't work now since
-// I have to link LLVMObject to do that (linker error otherwise).
-// Copy of llvm::object::OffloadKind.
-/// The producer of the associated offloading image.
-enum OffloadKind : uint16_t {
- OFK_None = 0,
- OFK_OpenMP = (1 << 0),
- OFK_Cuda = (1 << 1),
- OFK_HIP = (1 << 2),
- OFK_SYCL = (1 << 3),
- OFK_LAST = (1 << 4),
-};
-
-// Copy of llvm::object::ImageKind.
-/// The type of contents the offloading image contains.
-enum ImageKind : uint16_t {
- IMG_None = 0,
- IMG_Object,
- IMG_Bitcode,
- IMG_Cubin,
- IMG_Fatbinary,
- IMG_PTX,
- IMG_SPIRV,
- IMG_LAST,
-};
+static constexpr char DeviceBinaryTripleSPIRV64[] = "spirv64-unknown-unknown";
/// Device binary descriptor version supported by this library.
-static const uint16_t _LIBSYCL_SUPPORTED_DEVICE_BINARY_VERSION = 3;
+static constexpr uint16_t SupportedDevicyBinaryVersion = 3;
/// This struct is a record of the device binary information.
/// It must match the __tgt_device_image structure generated by the
-/// clang-offload-wrapper tool when their `Version` field match.
+/// clang-offload-wrapper tool when their `Version` fields match.
struct __sycl_tgt_device_image {
uint16_t Version;
/// The type of offload model the binary employs. See `OffloadKind`. Only
@@ -88,32 +47,32 @@ struct __sycl_tgt_device_image {
/// A null-terminated string representation of the device's target
/// architecture. Must hold one of _LIBSYCL_DEVICE_BINARY_TARGET_* values.
const char *TripleString;
- /// A null-terminated string; target- and compiler-specific options
- /// which are suggested to use to "compile" program at runtime.
+ /// A null-terminated string of target- and compiler-specific options
+ /// that are suggested to use to "compile" program at runtime.
const char *CompileOptions;
- /// A null-terminated string; target- and compiler-specific options
- /// which are suggested to use to "link" program at runtime.
+ /// A null-terminated string of target- and compiler-specific options
+ /// that are suggested to use to "link" program at runtime.
const char *LinkOptions;
/// Pointer to the target code start.
const unsigned char *ImageStart;
/// Pointer to the target code end.
const unsigned char *ImageEnd;
/// The offload entry table
- EntryTy *EntriesBegin;
- EntryTy *EntriesEnd;
+ llvm::offloading::EntryTy *EntriesBegin;
+ llvm::offloading::EntryTy *EntriesEnd;
// TODO: properties are not supported now.
- /// Array of preperty sets.
+ /// Array of property sets.
void *PropertiesBegin;
void *PropertiesEnd;
};
/// Version of offload binaries descriptor `__sycl_tgt_bin_desc` supported by
/// libsycl.
-static constexpr uint16_t _LIBSYCL_SUPPORTED_OFFLOAD_BINARY_VERSION = 1;
+static constexpr uint16_t SupportedOffloadBinaryVersion = 1;
/// This struct is a record of all the device code that may be offloaded.
/// It must match the `__tgt_bin_desc` structure generated by
-/// the clang-offload-wrapper tool when their `Version` field match.
+/// the clang-offload-wrapper tool when their `Version` fields match.
struct __sycl_tgt_bin_desc {
/// Version of the structure.
uint16_t Version;
@@ -122,8 +81,11 @@ struct __sycl_tgt_bin_desc {
/// Device binaries data.
__sycl_tgt_device_image *DeviceImages;
/// The offload entry table (not used, for compatibility with OpenMP).
- EntryTy *HostEntriesBegin;
- EntryTy *HostEntriesEnd;
+ llvm::offloading::EntryTy *HostEntriesBegin;
+ llvm::offloading::EntryTy *HostEntriesEnd;
};
+} // namespace detail
+_LIBSYCL_END_NAMESPACE_SYCL
+
#endif // _LIBSYCL_DEVICE_BINARY_STRUCTURES
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
index 534a4be4fcf1e..4b0be66eb97ae 100644
--- a/libsycl/src/detail/device_image_wrapper.hpp
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -16,9 +16,9 @@
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
-/// A wrapper of __sycl_tgt_device_image structure.
+/// A wrapper of __sycl_tgt_device_image structure to help with its fields
+/// parsing, iteration over data and data transformation.
class DeviceImageWrapper {
-public:
public:
DeviceImageWrapper(const __sycl_tgt_device_image &Bin) : MBin(&Bin) {}
// Explicitly delete copy constructor/operator= to avoid unintentional copies.
@@ -28,12 +28,13 @@ class DeviceImageWrapper {
DeviceImageWrapper(DeviceImageWrapper &&) = default;
DeviceImageWrapper &operator=(DeviceImageWrapper &&) = default;
- ~DeviceImageWrapper() {}
+ ~DeviceImageWrapper() = default;
- /// \return reference to corresponsing raw __sycl_tgt_device_image object.
+ /// \return a reference to the corresponding raw __sycl_tgt_device_image
+ /// object.
const __sycl_tgt_device_image &getRawData() const { return *get(); }
- /// \return size of corresponding device image data in bytes.
+ /// \return the size of the corresponding device image data in bytes.
size_t getSize() const {
return static_cast<size_t>(MBin->ImageEnd - MBin->ImageStart);
}
diff --git a/libsycl/src/detail/kernel_id.hpp b/libsycl/src/detail/kernel_id.hpp
index a1508042ed960..a8009dace290d 100644
--- a/libsycl/src/detail/kernel_id.hpp
+++ b/libsycl/src/detail/kernel_id.hpp
@@ -18,14 +18,14 @@
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
-/// The class is impl counterpart for sycl::kernel_id which represent a kernel
-/// identificator.
+/// The class is the implementation counterpart for sycl::kernel_id, which
+/// represents a kernel identificator.
class KernelIdImpl {
public:
- KernelIdImpl(std::string_view Name) : MName(std::move(Name)) {}
+ KernelIdImpl(std::string_view Name) : MName(Name) {}
KernelIdImpl() {}
- /// \return a null-terminated string representing the name of kernel this id
- /// stands for.
+ /// \return a null-terminated string representing the name of the kernel this
+ /// id stands for.
const char *get_name() { return MName.data(); }
private:
@@ -55,8 +55,8 @@ class kernel_id {
return !(lhs == rhs);
}
- /// \returns a null-terminated string which contains the kernel name.
- const char *get_name() const noexcept;
+ /// \returns a null-terminated string that contains the kernel name.
+ const char *get_name() const noexcept { return impl->get_name(); }
private:
kernel_id(const char *Name);
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index ea8673e0f39e9..f9b158a6e1918 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -6,11 +6,12 @@
//
//===----------------------------------------------------------------------===//
+#include <detail/program_manager.hpp>
+
#include <sycl/__impl/exception.hpp>
#include <detail/device_impl.hpp>
#include <detail/offload/offload_utils.hpp>
-#include <detail/program_manager.hpp>
#include <cstring>
@@ -18,14 +19,14 @@ _LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
static inline bool checkFatBinVersion(const __sycl_tgt_bin_desc &FatbinDesc) {
- return FatbinDesc.Version == _LIBSYCL_SUPPORTED_OFFLOAD_BINARY_VERSION;
+ return FatbinDesc.Version == SupportedOffloadBinaryVersion;
}
static inline bool
-checkDeviceImageValidness(const __sycl_tgt_device_image &DeviceImage) {
- return (DeviceImage.Version == _LIBSYCL_SUPPORTED_DEVICE_BINARY_VERSION) &&
- (DeviceImage.OffloadKind == OFK_SYCL) &&
- (DeviceImage.ImageFormat == IMG_SPIRV);
+checkDeviceImageValidity(const __sycl_tgt_device_image &DeviceImage) {
+ return (DeviceImage.Version == SupportedDevicyBinaryVersion) &&
+ (DeviceImage.OffloadKind == llvm::object::OFK_SYCL) &&
+ (DeviceImage.ImageFormat == llvm::object::IMG_SPIRV);
}
void ProgramManager::addImages(__sycl_tgt_bin_desc *FatbinDesc) {
@@ -38,14 +39,14 @@ void ProgramManager::addImages(__sycl_tgt_bin_desc *FatbinDesc) {
return;
std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
- for (int I = 0; I < FatbinDesc->NumDeviceBinaries; I++) {
+ for (int I = 0; I < FatbinDesc->NumDeviceBinaries; ++I) {
const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
- if (!checkDeviceImageValidness(RawDeviceImage))
+ if (!checkDeviceImageValidity(RawDeviceImage))
throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
"Incompatible device image.");
- const EntryTy *EntriesB = RawDeviceImage.EntriesBegin;
- const EntryTy *EntriesE = RawDeviceImage.EntriesEnd;
+ const llvm::offloading::EntryTy *EntriesB = RawDeviceImage.EntriesBegin;
+ const llvm::offloading::EntryTy *EntriesE = RawDeviceImage.EntriesEnd;
// Ignore "empty" device image.
if (EntriesB == EntriesE)
continue;
@@ -53,7 +54,7 @@ void ProgramManager::addImages(__sycl_tgt_bin_desc *FatbinDesc) {
std::unique_ptr<DeviceImageWrapper> NewImageWrapper =
std::make_unique<DeviceImageWrapper>(RawDeviceImage);
- for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt++) {
+ for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; ++EntriesIt) {
auto Name = EntriesIt->SymbolName;
auto KernelIDIt = MKernelNameToID.find(Name);
if (KernelIDIt == MKernelNameToID.end()) {
@@ -77,27 +78,24 @@ void ProgramManager::addImages(__sycl_tgt_bin_desc *FatbinDesc) {
void ProgramManager::removeImages(__sycl_tgt_bin_desc *FatbinDesc) {
assert(FatbinDesc && "Device images descriptor can't be nullptr");
- if (!checkFatBinVersion(*FatbinDesc))
- throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
- "Incompatible version of device images descriptor.");
- if (FatbinDesc->NumDeviceBinaries == 0)
+ if (!checkFatBinVersion(*FatbinDesc) || FatbinDesc->NumDeviceBinaries == 0)
return;
std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
- for (int I = 0; I < FatbinDesc->NumDeviceBinaries; I++) {
+ for (int I = 0; I < FatbinDesc->NumDeviceBinaries; ++I) {
const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
auto DevImageIt = MDeviceImageWrappers.find(&RawDeviceImage);
if (DevImageIt == MDeviceImageWrappers.end())
continue;
- const EntryTy *EntriesB = RawDeviceImage.EntriesBegin;
- const EntryTy *EntriesE = RawDeviceImage.EntriesEnd;
+ const llvm::offloading::EntryTy *EntriesB = RawDeviceImage.EntriesBegin;
+ const llvm::offloading::EntryTy *EntriesE = RawDeviceImage.EntriesEnd;
// Ignore "empty" device image
if (EntriesB == EntriesE)
continue;
- for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt++) {
+ for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; ++EntriesIt) {
if (auto KernelIDIt = MKernelNameToID.find(EntriesIt->SymbolName);
KernelIDIt != MKernelNameToID.end()) {
MKernelIDToDevImageJIT.erase(KernelIDIt->second);
@@ -114,18 +112,18 @@ static bool isImageTargetCompatible(const DeviceImageWrapper &Image,
sycl::backend BE = Device.getBackend();
const char *Target = Image.getRawData().TripleString;
- return (strcmp(Target, _LIBSYCL_DEVICE_BINARY_TARGET_SPIRV64) == 0) &&
+ return (strcmp(Target, DeviceBinaryTripleSPIRV64) == 0) &&
(BE == sycl::backend::level_zero);
}
-DeviceImageWrapper *ProgramManager::getDeviceImage(const char *KernelName,
- kernel_id KernelID,
+DeviceImageWrapper *ProgramManager::getDeviceImage(std::string_view KernelName,
+ const kernel_id &KernelID,
DeviceImpl &Device) {
+ std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
auto [Begin, End] = MKernelIDToDevImageJIT.equal_range(KernelID);
if (Begin != End) {
- ol_result_t Result{};
bool IsValid{};
- // TODO: with AOT (not implemented yet), we need to analize and check
+ // TODO: with AOT (not implemented yet), we need to analyze and check
// olIsValidBinary for AOT binaries first.
for (auto It = Begin; It != End; ++It) {
if (isImageTargetCompatible(*It->second, Device)) {
@@ -145,10 +143,12 @@ DeviceImageWrapper *ProgramManager::getDeviceImage(const char *KernelName,
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
-extern "C" void __sycl_register_lib(__sycl_tgt_bin_desc *FatbinDesc) {
+extern "C" _LIBSYCL_EXPORT void
+__sycl_register_lib(sycl::detail::__sycl_tgt_bin_desc *FatbinDesc) {
sycl::detail::ProgramManager::getInstance().addImages(FatbinDesc);
}
-extern "C" void __sycl_unregister_lib(__sycl_tgt_bin_desc *FatbinDesc) {
+extern "C" _LIBSYCL_EXPORT void
+__sycl_unregister_lib(sycl::detail::__sycl_tgt_bin_desc *FatbinDesc) {
sycl::detail::ProgramManager::getInstance().removeImages(FatbinDesc);
}
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index 3b1a2eb07d5ff..7d66602151d64 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -20,15 +20,16 @@
// +++ Entry points referenced by the offload wrapper object {
-/// Executed as a part of current module's (.exe, .dll) static initialization.
+/// Executed as a part of a module's (.exe, .dll) static initialization.
/// Registers device executable images with the runtime.
-extern "C" _LIBSYCL_EXPORT void __sycl_register_lib(__sycl_tgt_bin_desc *desc);
+extern "C" _LIBSYCL_EXPORT void
+__sycl_register_lib(sycl::detail::__sycl_tgt_bin_desc *FatbinDesc);
/// Executed as a part of current module's (.exe, .dll) static
/// de-initialization.
/// Unregisters device executable images with the runtime.
extern "C" _LIBSYCL_EXPORT void
-__sycl_unregister_lib(__sycl_tgt_bin_desc *desc);
+__sycl_unregister_lib(sycl::detail::__sycl_tgt_bin_desc *FatbinDesc);
// +++ }
@@ -47,17 +48,17 @@ class ProgramManager {
return PM;
}
- /// Parses raw device images data and prepare internal structures for
+ /// Parses raw device images data and prepares internal structures for
/// effective kernel/program creation.
- /// \param FatbinDesc a record of all the device code that may be offloaded
+ /// \param FatbinDesc a record of all the device code that may be offloaded,
/// generated by compiler and offloading tools.
- /// \throw sycl::exception with sycl::errc::runtime if device image descriptor
- /// has incompatible version or if device image has incompatible
- /// version/target/kind.
+ /// \throw sycl::exception with sycl::errc::runtime if a device image
+ /// descriptor has an incompatible version or if a device image has an
+ /// incompatible version, target or kind.
void addImages(__sycl_tgt_bin_desc *FatbinDesc);
- /// Removes all entries of the data in FatbinDesc in internal structures.
- /// \param FatbinDesc a record of all the device code that may be offloaded
+ /// Removes all entries of the data in FatbinDesc from internal structures.
+ /// \param FatbinDesc a record of all the device code that may be offloaded,
/// generated by compiler and offloading tools. Must match the pointer and
/// data passed to addImages.
void removeImages(__sycl_tgt_bin_desc *FatbinDesc);
@@ -68,15 +69,16 @@ class ProgramManager {
ProgramManager(ProgramManager const &) = delete;
ProgramManager &operator=(ProgramManager const &) = delete;
- /// Searches for device image that contains requested kernel and is compatible
- /// with requested device.
- /// \param KernelName a null-terminated string representing the name of kernel
- /// to obtain device image for.
+ /// Searches for a device image that contains the requested kernel and is
+ /// compatible with the requested device.
+ /// \param KernelName a null-terminated string representing the name of the
+ /// kernel to obtain a device image for.
/// \param KernelID a kernel id matching KernelName.
- /// \param DeviceImpl a device that device image must be compatible with.
- /// \throw sycl::exception with sycl::errc::runtime if device image validness
- /// check failed in liboffload or if no compatible image was found.
- DeviceImageWrapper *getDeviceImage(const char *KernelName, kernel_id KernelID,
+ /// \param DeviceImpl a device with which device image must be compatible.
+ /// \throw sycl::exception with sycl::errc::runtime if the device image
+ /// validation failed in liboffload or if no compatible image was found.
+ DeviceImageWrapper *getDeviceImage(std::string_view KernelName,
+ const kernel_id &KernelID,
DeviceImpl &Device);
// Filled by addImages(...).
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index dec2d7d5507aa..9c93fe02de8a6 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -6,9 +6,10 @@
//
//===----------------------------------------------------------------------===//
-#include <detail/device_impl.hpp>
#include <detail/queue_impl.hpp>
+#include <detail/device_impl.hpp>
+
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
>From aaa76ae7598ee6af2aabdeb4dd7931c5f250cb5a Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 25 Mar 2026 02:50:37 -0700
Subject: [PATCH 3/9] remove clang-offload-wrapper mentioning
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/detail/device_binary_structures.hpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/libsycl/src/detail/device_binary_structures.hpp b/libsycl/src/detail/device_binary_structures.hpp
index be403da9fcbb4..3b94ddca1f2d8 100644
--- a/libsycl/src/detail/device_binary_structures.hpp
+++ b/libsycl/src/detail/device_binary_structures.hpp
@@ -36,7 +36,7 @@ static constexpr uint16_t SupportedDevicyBinaryVersion = 3;
/// This struct is a record of the device binary information.
/// It must match the __tgt_device_image structure generated by the
-/// clang-offload-wrapper tool when their `Version` fields match.
+/// compiler when their `Version` fields match.
struct __sycl_tgt_device_image {
uint16_t Version;
/// The type of offload model the binary employs. See `OffloadKind`. Only
@@ -72,7 +72,7 @@ static constexpr uint16_t SupportedOffloadBinaryVersion = 1;
/// This struct is a record of all the device code that may be offloaded.
/// It must match the `__tgt_bin_desc` structure generated by
-/// the clang-offload-wrapper tool when their `Version` fields match.
+/// the compiler when their `Version` fields match.
struct __sycl_tgt_bin_desc {
/// Version of the structure.
uint16_t Version;
>From 3683423bfa6e1de9fa2fba8bee44285d81b8587c Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 25 Mar 2026 04:02:52 -0700
Subject: [PATCH 4/9] [libsycl] add sycl::event and wait functionality to event
& queue
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/docs/index.rst | 11 +++-
libsycl/include/sycl/__impl/event.hpp | 90 +++++++++++++++++++++++++++
libsycl/include/sycl/__impl/queue.hpp | 7 +++
libsycl/include/sycl/sycl.hpp | 1 +
libsycl/src/CMakeLists.txt | 2 +
libsycl/src/detail/event_impl.cpp | 39 ++++++++++++
libsycl/src/detail/event_impl.hpp | 68 ++++++++++++++++++++
libsycl/src/detail/queue_impl.cpp | 14 ++++-
libsycl/src/detail/queue_impl.hpp | 9 ++-
libsycl/src/event.cpp | 25 ++++++++
libsycl/src/queue.cpp | 2 +
11 files changed, 264 insertions(+), 4 deletions(-)
create mode 100644 libsycl/include/sycl/__impl/event.hpp
create mode 100644 libsycl/src/detail/event_impl.cpp
create mode 100644 libsycl/src/detail/event_impl.hpp
create mode 100644 libsycl/src/event.cpp
diff --git a/libsycl/docs/index.rst b/libsycl/docs/index.rst
index 03f7fb7c0876e..9aa36b4a54c57 100644
--- a/libsycl/docs/index.rst
+++ b/libsycl/docs/index.rst
@@ -106,7 +106,14 @@ TODO for added SYCL classes
* device selection: to add compatibility with old SYCL 1.2.1 device selectors, still part of SYCL 2020 specification
* ``context``: to implement get_info, properties & public constructors once context support is added to liboffload
-* ``queue``: to implement USM methods, to implement synchronization methods, to implement submit & copy with accessors (low priority), get_info & properties, ctors that accepts context (blocked by lack of liboffload support)
+* ``queue``:
+
+ * to implement USM methods
+ * to implement synchronization methods
+ * to implement submit & copy with accessors (low priority)
+ * get_info & properties
+ * ctors that accepts context (blocked by lack of liboffload support)
+
* ``property_list``: to fully implement and integrate with existing SYCL runtime classes supporting it
* usm allocations:
@@ -114,7 +121,7 @@ TODO for added SYCL classes
* forward templated funcs to alignment methods (rewrite current impl)
* handle sub devices once they are implemented (blocked by liboffload support)
-
+* ``event``: get_wait_list, get_info, get_profiling_info, wait_and_throw & default ctor are not implemented
* general opens:
* define a way to report errors from object dtors.
\ No newline at end of file
diff --git a/libsycl/include/sycl/__impl/event.hpp b/libsycl/include/sycl/__impl/event.hpp
new file mode 100644
index 0000000000000..7df095c9a1fd1
--- /dev/null
+++ b/libsycl/include/sycl/__impl/event.hpp
@@ -0,0 +1,90 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+/// This file contains the declaration of the SYCL event class (SYCL
+/// 2020 4.6.6.), that represents the status of an operation that is being
+/// executed by the SYCL runtime.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_EVENT_HPP
+#define _LIBSYCL___IMPL_EVENT_HPP
+
+#include <sycl/__impl/backend.hpp>
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/detail/obj_utils.hpp>
+#include <sycl/__impl/info/desc_base.hpp>
+
+#include <memory>
+#include <vector>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+class event;
+
+namespace detail {
+class EventImpl;
+template <typename T>
+using is_event_info_desc_t = typename is_info_desc<T, event>::return_type;
+} // namespace detail
+
+/// SYCL 2020 4.6.6. Event class.
+class _LIBSYCL_EXPORT event {
+public:
+ event(const event &rhs) = default;
+
+ event(event &&rhs) = default;
+
+ event &operator=(const event &rhs) = default;
+
+ event &operator=(event &&rhs) = default;
+
+ friend bool operator==(const event &lhs, const event &rhs) {
+ return lhs.impl == rhs.impl;
+ }
+
+ friend bool operator!=(const event &lhs, const event &rhs) {
+ return !(lhs == rhs);
+ }
+
+ /// \return the backend associated with this platform.
+ backend get_backend() const noexcept;
+
+ /// Blocks until all commands associated with this event and any dependent
+ /// events have completed.
+ void wait();
+
+ /// Behaves as if calling event::wait on each event in eventList.
+ static void wait(const std::vector<event> &eventList);
+
+ /// Queries this SYCL event for information.
+ ///
+ /// \return depends on the information being requested.
+ template <typename Param>
+ detail::is_event_info_desc_t<Param> get_info() const;
+
+ /// Queries this SYCL event for SYCL backend-specific information.
+ ///
+ /// \return depends on information being queried.
+ template <typename Param>
+ typename Param::return_type get_backend_info() const;
+
+private:
+ event(std::shared_ptr<detail::EventImpl> Impl) : impl(Impl) {}
+ std::shared_ptr<detail::EventImpl> impl;
+
+ friend sycl::detail::ImplUtils;
+};
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+template <>
+struct std::hash<sycl::event> : public sycl::detail::HashBase<sycl::event> {};
+
+#endif // _LIBSYCL___IMPL_EVENT_HPP
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index a440959c6311f..587f56a8eb245 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -17,6 +17,7 @@
#include <sycl/__impl/async_handler.hpp>
#include <sycl/__impl/device.hpp>
+#include <sycl/__impl/event.hpp>
#include <sycl/__impl/property_list.hpp>
#include <sycl/__impl/detail/config.hpp>
@@ -29,6 +30,7 @@ class context;
namespace detail {
class QueueImpl;
+
} // namespace detail
// SYCL 2020 4.6.5. Queue class.
@@ -136,6 +138,11 @@ class _LIBSYCL_EXPORT queue {
template <typename Param>
typename Param::return_type get_backend_info() const;
+ /// Blocks the calling thread until all commands previously submitted to this
+ /// queue have completed. Synchronous errors are reported through SYCL
+ /// exceptions.
+ void wait();
+
private:
queue(const std::shared_ptr<detail::QueueImpl> &Impl) : impl(Impl) {}
std::shared_ptr<detail::QueueImpl> impl;
diff --git a/libsycl/include/sycl/sycl.hpp b/libsycl/include/sycl/sycl.hpp
index 3fcf088f45535..ce9fc8defd90b 100644
--- a/libsycl/include/sycl/sycl.hpp
+++ b/libsycl/include/sycl/sycl.hpp
@@ -17,6 +17,7 @@
#include <sycl/__impl/context.hpp>
#include <sycl/__impl/device.hpp>
#include <sycl/__impl/device_selector.hpp>
+#include <sycl/__impl/event.hpp>
#include <sycl/__impl/exception.hpp>
#include <sycl/__impl/platform.hpp>
#include <sycl/__impl/queue.hpp>
diff --git a/libsycl/src/CMakeLists.txt b/libsycl/src/CMakeLists.txt
index 4501005e433e3..7b9826fb8a3de 100644
--- a/libsycl/src/CMakeLists.txt
+++ b/libsycl/src/CMakeLists.txt
@@ -85,6 +85,7 @@ endfunction(add_sycl_rt_library)
set(LIBSYCL_SOURCES
"context.cpp"
+ "event.cpp"
"exception.cpp"
"exception_list.cpp"
"device.cpp"
@@ -93,6 +94,7 @@ set(LIBSYCL_SOURCES
"queue.cpp"
"usm_functions.cpp"
"detail/context_impl.cpp"
+ "detail/event_impl.cpp"
"detail/device_impl.cpp"
"detail/global_objects.cpp"
"detail/platform_impl.cpp"
diff --git a/libsycl/src/detail/event_impl.cpp b/libsycl/src/detail/event_impl.cpp
new file mode 100644
index 0000000000000..895f8029d4c35
--- /dev/null
+++ b/libsycl/src/detail/event_impl.cpp
@@ -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
+//
+//===----------------------------------------------------------------------===//
+
+#include <detail/event_impl.hpp>
+#include <detail/platform_impl.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+EventImpl::~EventImpl() {
+ if (MOffloadEvent)
+ std::ignore = olDestroyEvent(MOffloadEvent);
+}
+
+backend EventImpl::getBackend() const noexcept {
+ // TODO: to handle default cosntructed
+ // The event is constructed as though it were created from a
+ // default-constructed queue. Therefore, its backend is the same as the
+ // backend of the device selected by default_selector_v.
+ return MPlatform.getBackend();
+}
+
+void EventImpl::wait() {
+ // MOffloadEvent == nullptr when event is default constructed. Default
+ // constructed event is immediately ready.
+ if (!MOffloadEvent)
+ return;
+
+ callAndThrow(olSyncEvent, MOffloadEvent);
+}
+
+} // namespace detail
+_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/event_impl.hpp b/libsycl/src/detail/event_impl.hpp
new file mode 100644
index 0000000000000..f570538512def
--- /dev/null
+++ b/libsycl/src/detail/event_impl.hpp
@@ -0,0 +1,68 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_EVENT_IMPL
+#define _LIBSYCL_EVENT_IMPL
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/queue.hpp>
+
+#include <OffloadAPI.h>
+
+#include <memory>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+namespace detail {
+
+class PlatformImpl;
+
+class EventImpl {
+ // Helper to limit EventImpl creation.
+ struct PrivateTag {
+ explicit PrivateTag() = default;
+ };
+
+public:
+ /// Constructs a SYCL event instance using the provided
+ /// offload event instance.
+ ///
+ /// \param Event is a raw offload library handle representing event.
+ /// \param Platform is a platform this event belongs to.
+ EventImpl(ol_event_handle_t Event, PlatformImpl &Platform, PrivateTag)
+ : MOffloadEvent(Event), MPlatform(Platform) {}
+
+ static std::shared_ptr<EventImpl>
+ createEventWithHandle(ol_event_handle_t Event, PlatformImpl &Queue) {
+ return std::make_shared<EventImpl>(Event, Queue, PrivateTag{});
+ }
+
+ /// Releases handle to the corresponding liboffload event.
+ ~EventImpl();
+
+ /// \return the sycl::backend associated with this event.
+ backend getBackend() const noexcept;
+
+ /// Waits for completion of the corresponding kernel and its dependencies.
+ void wait();
+
+ /// \return liboffload handle that this SYCL event represents.
+ ol_event_handle_t getHandle() { return MOffloadEvent; }
+
+ /// \return a platform implementation object this event belongs to.
+ const PlatformImpl &getPlatformImpl() const { return MPlatform; }
+
+private:
+ ol_event_handle_t MOffloadEvent{};
+ PlatformImpl &MPlatform;
+};
+
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL_EVENT_IMPL
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index 9c93fe02de8a6..74ccc48877c09 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -9,6 +9,8 @@
#include <detail/queue_impl.hpp>
#include <detail/device_impl.hpp>
+#include <detail/event_impl.hpp>
+#include <detail/program_manager.hpp>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
@@ -18,9 +20,19 @@ QueueImpl::QueueImpl(DeviceImpl &deviceImpl, const async_handler &asyncHandler,
const property_list &propList, PrivateTag)
: MIsInorder(false), MAsyncHandler(asyncHandler), MPropList(propList),
MDevice(deviceImpl),
- MContext(MDevice.getPlatformImpl().getDefaultContext()) {}
+ MContext(MDevice.getPlatformImpl().getDefaultContext()) {
+ callAndThrow(olCreateQueue, MDevice.getHandle(), &MOffloadQueue);
+}
+
+QueueImpl::~QueueImpl() {
+ // TODO: consider where to report errors
+ if (MOffloadQueue)
+ std::ignore = olDestroyQueue(MOffloadQueue);
+}
backend QueueImpl::getBackend() const noexcept { return MDevice.getBackend(); }
+void QueueImpl::wait() { callAndThrow(olSyncQueue, MOffloadQueue); }
+
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/queue_impl.hpp b/libsycl/src/detail/queue_impl.hpp
index 6403099a19060..cdb7595e852ec 100644
--- a/libsycl/src/detail/queue_impl.hpp
+++ b/libsycl/src/detail/queue_impl.hpp
@@ -21,6 +21,9 @@ namespace detail {
class ContextImpl;
class DeviceImpl;
+class EventImpl;
+
+using EventImplPtr = std::shared_ptr<EventImpl>;
class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
struct PrivateTag {
@@ -28,7 +31,7 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
};
public:
- ~QueueImpl() = default;
+ ~QueueImpl();
/// Constructs a SYCL queue from a device using an asyncHandler and
/// a propList.
@@ -59,7 +62,11 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
/// \return true if and only if the queue is in order.
bool isInOrder() const { return MIsInorder; }
+ /// Waits for completion of all kernels submitted to this queue.
+ void wait();
+
private:
+ ol_queue_handle_t MOffloadQueue = {};
const bool MIsInorder;
const async_handler MAsyncHandler;
const property_list MPropList;
diff --git a/libsycl/src/event.cpp b/libsycl/src/event.cpp
new file mode 100644
index 0000000000000..68046211272f2
--- /dev/null
+++ b/libsycl/src/event.cpp
@@ -0,0 +1,25 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <sycl/__impl/event.hpp>
+
+#include <detail/event_impl.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+backend event::get_backend() const noexcept { return impl->getBackend(); }
+
+void event::wait(const std::vector<event> &EventList) {
+ for (auto Event : EventList) {
+ Event.wait();
+ }
+}
+
+void event::wait() { return impl->wait(); }
+
+_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp
index faed274674447..9fe020eabf2cc 100644
--- a/libsycl/src/queue.cpp
+++ b/libsycl/src/queue.cpp
@@ -33,4 +33,6 @@ device queue::get_device() const {
bool queue::is_in_order() const { return impl->isInOrder(); }
+void queue::wait() { return impl->wait(); }
+
_LIBSYCL_END_NAMESPACE_SYCL
>From 1394cc9550ce4b875b84b3e622cc4fb84fef8155 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 25 Mar 2026 04:56:49 -0700
Subject: [PATCH 5/9] [libsycl] Add kernel creation
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/detail/program_manager.cpp | 100 ++++++++++++++++++++++++-
libsycl/src/detail/program_manager.hpp | 87 ++++++++++++++++++++-
2 files changed, 184 insertions(+), 3 deletions(-)
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index f9b158a6e1918..7d6523daeb6ee 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -18,6 +18,20 @@
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
+ProgramWrapper::ProgramWrapper(ol_device_handle_t Device,
+ DeviceImageWrapper &DevImage) {
+ assert(Device);
+
+ callAndThrow(olCreateProgram, Device, DevImage.getRawData().ImageStart,
+ DevImage.getSize(), &MProgram);
+}
+
+ProgramWrapper::~ProgramWrapper() {
+ assert(MProgram);
+ std::ignore = olDestroyProgram(MProgram);
+ // TODO: define a way to report errors from dtors.
+}
+
static inline bool checkFatBinVersion(const __sycl_tgt_bin_desc &FatbinDesc) {
return FatbinDesc.Version == SupportedOffloadBinaryVersion;
}
@@ -81,7 +95,7 @@ void ProgramManager::removeImages(__sycl_tgt_bin_desc *FatbinDesc) {
if (!checkFatBinVersion(*FatbinDesc) || FatbinDesc->NumDeviceBinaries == 0)
return;
- std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
+ std::scoped_lock Guard{MImageCollectionMutex, MKernelCollectionsMutex};
for (int I = 0; I < FatbinDesc->NumDeviceBinaries; ++I) {
const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
@@ -99,10 +113,18 @@ void ProgramManager::removeImages(__sycl_tgt_bin_desc *FatbinDesc) {
if (auto KernelIDIt = MKernelNameToID.find(EntriesIt->SymbolName);
KernelIDIt != MKernelNameToID.end()) {
MKernelIDToDevImageJIT.erase(KernelIDIt->second);
+ MKernels.erase(KernelIDIt->second);
MKernelNameToID.erase(KernelIDIt);
}
}
+ if (auto ProgramIt = MPrograms.find(DevImageIt->second.get());
+ ProgramIt != MPrograms.end()) {
+ for (auto &[Device, Program] : ProgramIt->second) {
+ MProgramWrappers.erase(Program);
+ MPrograms.erase(ProgramIt);
+ }
+ }
MDeviceImageWrappers.erase(DevImageIt);
}
}
@@ -119,7 +141,6 @@ static bool isImageTargetCompatible(const DeviceImageWrapper &Image,
DeviceImageWrapper *ProgramManager::getDeviceImage(std::string_view KernelName,
const kernel_id &KernelID,
DeviceImpl &Device) {
- std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
auto [Begin, End] = MKernelIDToDevImageJIT.equal_range(KernelID);
if (Begin != End) {
bool IsValid{};
@@ -140,6 +161,81 @@ DeviceImageWrapper *ProgramManager::getDeviceImage(std::string_view KernelName,
"No kernel named " + std::string(KernelName) + " was found");
}
+ol_symbol_handle_t ProgramManager::getOrCreateKernel(const char *KernelName,
+ DeviceImpl &Device) {
+ std::lock_guard<std::mutex> ImageGuard(MImageCollectionMutex);
+
+ auto KernelIDIt = MKernelNameToID.find(KernelName);
+ if (KernelIDIt == MKernelNameToID.end())
+ throw exception(make_error_code(errc::runtime),
+ "No kernel named " + std::string(KernelName) +
+ " was found");
+
+ std::lock_guard<std::mutex> KernelGuard(MKernelCollectionsMutex);
+
+ auto Kernel = getKernel(KernelIDIt->second, Device);
+ if (Kernel)
+ return Kernel;
+
+ DeviceImageWrapper *DevImage =
+ getDeviceImage(KernelName, KernelIDIt->second, Device);
+ if (!DevImage)
+ throw;
+
+ ol_program_handle_t Program = getOrCreateProgram(Device, DevImage);
+ assert(Program);
+ Kernel = createKernel(Program, KernelIDIt->second, KernelName, Device);
+ assert(Kernel);
+ return Kernel;
+}
+
+ol_program_handle_t
+ProgramManager::getOrCreateProgram(DeviceImpl &Device,
+ DeviceImageWrapper *DevImage) {
+ if (auto DevToProgramIt = MPrograms.find(DevImage);
+ DevToProgramIt != MPrograms.end()) {
+ auto ProgramIt = DevToProgramIt->second.find(Device.getHandle());
+ if (ProgramIt != DevToProgramIt->second.end())
+ return ProgramIt->second;
+ }
+
+ std::unique_ptr<ProgramWrapper> NewProgramWrapper(
+ new ProgramWrapper(Device.getHandle(), *DevImage));
+ auto Program = NewProgramWrapper->getHandle();
+ {
+ MPrograms[DevImage].insert(std::make_pair(Device.getHandle(), Program));
+ MProgramWrappers.insert(std::make_pair(NewProgramWrapper->getHandle(),
+ std::move(NewProgramWrapper)));
+ }
+
+ return Program;
+}
+
+ol_symbol_handle_t ProgramManager::createKernel(ol_program_handle_t Program,
+ const kernel_id &KernelID,
+ const char *KernelName,
+ DeviceImpl &Device) {
+ ol_symbol_handle_t Kernel{};
+ callAndThrow(olGetSymbol, Program, KernelName, OL_SYMBOL_KIND_KERNEL,
+ &Kernel);
+ MKernels.insert(
+ std::make_pair(KernelID, std::make_pair(Device.getHandle(), Kernel)));
+ return Kernel;
+}
+
+ol_symbol_handle_t ProgramManager::getKernel(const kernel_id &KernelID,
+ DeviceImpl &Device) {
+ auto Range = MKernels.equal_range(KernelID);
+ for (auto Kernels = Range.first; Kernels != Range.second; ++Kernels) {
+ auto &[KernelDevice, KernelSymbol] = Kernels->second;
+ if (KernelDevice == Device.getHandle()) {
+ assert(KernelSymbol && "Built kernel symbol can't be null");
+ return KernelSymbol;
+ }
+ }
+ return nullptr;
+}
+
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index 7d66602151d64..b017383a16b4c 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -15,6 +15,8 @@
#include <detail/device_image_wrapper.hpp>
#include <detail/kernel_id.hpp>
+#include <OffloadAPI.h>
+
#include <mutex>
#include <unordered_map>
@@ -39,6 +41,30 @@ namespace detail {
class DeviceImpl;
+/// A wrapper of liboffload program handle to manage its lifetime.
+class ProgramWrapper {
+public:
+ /// Constructs ProgramWrapper by creating liboffload program with the provided
+ /// arguments.
+ ///
+ /// \param Device is a device to use for program creation.
+ /// \param DevImage is a device image (wrapped __sycl_tgt_device_image) to use
+ /// for program creation.
+ /// \throw sycl::exception with sycl::errc::runtime when failed to create
+ /// program.
+ ProgramWrapper(ol_device_handle_t Device, DeviceImageWrapper &DevImage);
+
+ /// Releases the corresponding liboffload program handle by calling
+ /// olDestroyProgram.
+ ~ProgramWrapper();
+
+ /// \return the corresponding liboffload program handle.
+ ol_program_handle_t getHandle() { return MProgram; }
+
+private:
+ ol_program_handle_t MProgram{};
+};
+
/// A class to manage programs and kernels.
class ProgramManager {
@@ -63,6 +89,16 @@ class ProgramManager {
/// data passed to addImages.
void removeImages(__sycl_tgt_bin_desc *FatbinDesc);
+ /// Creates liboffload kernel that is ready for execution.
+ /// Thread-safe.
+ /// \param KernelName a null-terminated string representing a name of kernel
+ /// to be created.
+ /// \param Device a device for which this kernel must be compiled.
+ /// \return liboffload kernel handle that is ready to be passed to kernel
+ /// execution methods.
+ ol_symbol_handle_t getOrCreateKernel(const char *KernelName,
+ DeviceImpl &Device);
+
private:
ProgramManager() = default;
~ProgramManager() = default;
@@ -74,13 +110,43 @@ class ProgramManager {
/// \param KernelName a null-terminated string representing the name of the
/// kernel to obtain a device image for.
/// \param KernelID a kernel id matching KernelName.
- /// \param DeviceImpl a device with which device image must be compatible.
+ /// \param Device a device with which device image must be compatible.
/// \throw sycl::exception with sycl::errc::runtime if the device image
/// validation failed in liboffload or if no compatible image was found.
DeviceImageWrapper *getDeviceImage(std::string_view KernelName,
const kernel_id &KernelID,
DeviceImpl &Device);
+ /// Searches for or creates a program.
+ /// This call must be protected with mutex since it updates MPrograms and
+ /// MProgramWrappers collections.
+ /// \param Device a device that program must be created with.
+ /// \param DevImage a device image to get or create program with.
+ /// \return liboffload program for the requested configuration.
+ ol_program_handle_t getOrCreateProgram(DeviceImpl &Device,
+ DeviceImageWrapper *DevImage);
+
+ /// Creates kernel from program.
+ /// This call must be protected with mutex since it updates MKernels
+ /// collection.
+ /// \param Program a program to create kernel with.
+ /// \param KernelID an id of kernel to create.
+ /// \param KernelName a null-terminated string representing the name of kernel
+ /// to create.
+ /// \param Device a device that kernel must be created with.
+ /// \return liboffload kernel for the requested configuration.
+ ol_symbol_handle_t createKernel(ol_program_handle_t Program,
+ const kernel_id &KernelID,
+ const char *KernelName, DeviceImpl &Device);
+
+ /// Searches for kernel.
+ /// This call must be protected with mutex since it reads MKernels collection.
+ /// \param KernelID an id of kernel to look for.
+ /// \param Device a device that kernel must be created with.
+ /// \return liboffload kernel for the requested configuration or nullptr if
+ /// such kernel is not found.
+ ol_symbol_handle_t getKernel(const kernel_id &KernelID, DeviceImpl &Device);
+
// Filled by addImages(...).
std::unordered_map<std::string_view, kernel_id> MKernelNameToID;
std::unordered_map<kernel_id, DeviceImageWrapper *> MKernelIDToDevImageJIT;
@@ -88,7 +154,26 @@ class ProgramManager {
std::unordered_map<const __sycl_tgt_device_image *,
std::unique_ptr<DeviceImageWrapper>>
MDeviceImageWrappers;
+ // All data collections, created from data in __sycl_register_lib, must be
+ // protected with this mutex. Protects data that can be modified only by
+ // modules load/unload.
std::mutex MImageCollectionMutex;
+
+ // Filled by getOrCreateKernel and everything it calls inside.
+ std::unordered_map<
+ DeviceImageWrapper *,
+ std::unordered_map<ol_device_handle_t, ol_program_handle_t>>
+ MPrograms;
+ std::unordered_multimap<kernel_id,
+ std::pair<ol_device_handle_t, ol_symbol_handle_t>>
+ MKernels;
+
+ // Controls lifetime of programs.
+ std::unordered_map<ol_program_handle_t, std::unique_ptr<ProgramWrapper>>
+ MProgramWrappers;
+ // All data collections, used and modified by kernel submissions, must be
+ // protected with this mutex.
+ std::mutex MKernelCollectionsMutex;
};
} // namespace detail
>From 67a4c90a5b912b25107e0b49e88785033f997d57 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 1 Apr 2026 03:36:56 -0700
Subject: [PATCH 6/9] fix comments
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/detail/program_manager.cpp | 18 +++++++++---------
libsycl/src/detail/program_manager.hpp | 5 +++++
2 files changed, 14 insertions(+), 9 deletions(-)
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index 7d6523daeb6ee..18b892d272432 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -122,8 +122,8 @@ void ProgramManager::removeImages(__sycl_tgt_bin_desc *FatbinDesc) {
ProgramIt != MPrograms.end()) {
for (auto &[Device, Program] : ProgramIt->second) {
MProgramWrappers.erase(Program);
- MPrograms.erase(ProgramIt);
}
+ MPrograms.erase(ProgramIt);
}
MDeviceImageWrappers.erase(DevImageIt);
}
@@ -163,7 +163,7 @@ DeviceImageWrapper *ProgramManager::getDeviceImage(std::string_view KernelName,
ol_symbol_handle_t ProgramManager::getOrCreateKernel(const char *KernelName,
DeviceImpl &Device) {
- std::lock_guard<std::mutex> ImageGuard(MImageCollectionMutex);
+ std::unique_lock<std::mutex> ImageGuard(MImageCollectionMutex);
auto KernelIDIt = MKernelNameToID.find(KernelName);
if (KernelIDIt == MKernelNameToID.end())
@@ -179,8 +179,8 @@ ol_symbol_handle_t ProgramManager::getOrCreateKernel(const char *KernelName,
DeviceImageWrapper *DevImage =
getDeviceImage(KernelName, KernelIDIt->second, Device);
- if (!DevImage)
- throw;
+
+ ImageGuard.unlock();
ol_program_handle_t Program = getOrCreateProgram(Device, DevImage);
assert(Program);
@@ -202,11 +202,9 @@ ProgramManager::getOrCreateProgram(DeviceImpl &Device,
std::unique_ptr<ProgramWrapper> NewProgramWrapper(
new ProgramWrapper(Device.getHandle(), *DevImage));
auto Program = NewProgramWrapper->getHandle();
- {
- MPrograms[DevImage].insert(std::make_pair(Device.getHandle(), Program));
- MProgramWrappers.insert(std::make_pair(NewProgramWrapper->getHandle(),
- std::move(NewProgramWrapper)));
- }
+ MPrograms[DevImage].insert(std::make_pair(Device.getHandle(), Program));
+ MProgramWrappers.insert(std::make_pair(NewProgramWrapper->getHandle(),
+ std::move(NewProgramWrapper)));
return Program;
}
@@ -215,6 +213,8 @@ ol_symbol_handle_t ProgramManager::createKernel(ol_program_handle_t Program,
const kernel_id &KernelID,
const char *KernelName,
DeviceImpl &Device) {
+ assert((getKernel(KernelID, Device) == nullptr) &&
+ "Attempt to create kernel that already exists.");
ol_symbol_handle_t Kernel{};
callAndThrow(olGetSymbol, Program, KernelName, OL_SYMBOL_KIND_KERNEL,
&Kernel);
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index b017383a16b4c..22e380334befd 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -58,6 +58,11 @@ class ProgramWrapper {
/// olDestroyProgram.
~ProgramWrapper();
+ ProgramWrapper(const ProgramWrapper &) = delete;
+ ProgramWrapper &operator=(const ProgramWrapper &) = delete;
+ ProgramWrapper(ProgramWrapper &&) = delete;
+ ProgramWrapper &operator=(ProgramWrapper &&) = delete;
+
/// \return the corresponding liboffload program handle.
ol_program_handle_t getHandle() { return MProgram; }
>From 470543acf33600f82041bfb688d1efc77b6ade7d Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Fri, 10 Apr 2026 04:28:54 -0700
Subject: [PATCH 7/9] remove extra mutex
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/include/sycl/__impl/queue.hpp | 2 --
libsycl/src/detail/program_manager.cpp | 13 ++++++-------
libsycl/src/detail/program_manager.hpp | 20 ++++++++------------
3 files changed, 14 insertions(+), 21 deletions(-)
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 36f0b9ad57d4a..41b018b681b8e 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -17,7 +17,6 @@
#include <sycl/__impl/async_handler.hpp>
#include <sycl/__impl/device.hpp>
-#include <sycl/__impl/event.hpp>
#include <sycl/__impl/property_list.hpp>
#include <sycl/__impl/detail/config.hpp>
@@ -30,7 +29,6 @@ class context;
namespace detail {
class QueueImpl;
-
} // namespace detail
// SYCL 2020 4.6.5. Queue class.
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index aacd33b405806..f461c0bcdfec1 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -52,7 +52,7 @@ void ProgramAndKernelManager::registerFatBin(__sycl_tgt_bin_desc *FatbinDesc) {
if (!FatbinDesc->NumDeviceBinaries)
return;
- std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
+ std::lock_guard<std::mutex> Guard(MDataCollectionMutex);
for (uint16_t I = 0; I < FatbinDesc->NumDeviceBinaries; ++I) {
const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
if (!checkDeviceImageValidity(RawDeviceImage))
@@ -96,7 +96,7 @@ void ProgramAndKernelManager::unregisterFatBin(
if (!checkFatBinVersion(*FatbinDesc) || FatbinDesc->NumDeviceBinaries == 0)
return;
- std::scoped_lock Guard{MImageCollectionMutex, MKernelCollectionsMutex};
+ std::lock_guard<std::mutex> Guard(MDataCollectionMutex);
for (uint16_t I = 0; I < FatbinDesc->NumDeviceBinaries; ++I) {
const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
@@ -144,7 +144,7 @@ DeviceImageWrapper *
ProgramAndKernelManager::getDeviceImage(std::string_view KernelName,
const kernel_id &KernelID,
DeviceImpl &Device) {
- std::lock_guard<std::mutex> Guard(MImageCollectionMutex);
+ std::lock_guard<std::mutex> Guard(MDataCollectionMutex);
auto [Begin, End] = MKernelIDToDevImageJIT.equal_range(KernelID);
if (Begin != End) {
bool IsValid{};
@@ -167,7 +167,7 @@ ProgramAndKernelManager::getDeviceImage(std::string_view KernelName,
ol_symbol_handle_t ProgramAndKernelManager::getOrCreateKernel(const char *KernelName,
DeviceImpl &Device) {
- std::unique_lock<std::mutex> ImageGuard(MImageCollectionMutex);
+ std::unique_lock<std::mutex> ImageGuard(MDataCollectionMutex);
auto KernelIDIt = MKernelNameToID.find(KernelName);
if (KernelIDIt == MKernelNameToID.end())
@@ -175,7 +175,7 @@ ol_symbol_handle_t ProgramAndKernelManager::getOrCreateKernel(const char *Kernel
"No kernel named " + std::string(KernelName) +
" was found");
- std::lock_guard<std::mutex> KernelGuard(MKernelCollectionsMutex);
+ std::lock_guard<std::mutex> KernelGuard(MDataCollectionMutex);
auto Kernel = getKernel(KernelIDIt->second, Device);
if (Kernel)
@@ -184,10 +184,9 @@ ol_symbol_handle_t ProgramAndKernelManager::getOrCreateKernel(const char *Kernel
DeviceImageWrapper *DevImage =
getDeviceImage(KernelName, KernelIDIt->second, Device);
- ImageGuard.unlock();
-
ol_program_handle_t Program = getOrCreateProgram(Device, DevImage);
assert(Program);
+
Kernel = createKernel(Program, KernelIDIt->second, KernelName, Device);
assert(Kernel);
return Kernel;
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index bcf8be30f2407..5ab828b25a470 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -118,6 +118,7 @@ class ProgramAndKernelManager {
/// Searches for a device image that contains the requested kernel and is
/// compatible with the requested device.
+ /// This call must be protected with MDataCollectionMutex.
/// \param KernelName a null-terminated string representing the name of the
/// kernel to obtain a device image for.
/// \param KernelID a kernel id matching KernelName.
@@ -129,8 +130,8 @@ class ProgramAndKernelManager {
DeviceImpl &Device);
/// Searches for or creates a program.
- /// This call must be protected with mutex since it updates MPrograms and
- /// MProgramWrappers collections.
+ /// This call must be protected with MDataCollectionMutex since it updates
+ /// MPrograms and MProgramWrappers collections.
/// \param Device a device that program must be created with.
/// \param DevImage a device image to get or create program with.
/// \return liboffload program for the requested configuration.
@@ -138,8 +139,8 @@ class ProgramAndKernelManager {
DeviceImageWrapper *DevImage);
/// Creates kernel from program.
- /// This call must be protected with mutex since it updates MKernels
- /// collection.
+ /// This call must be protected with MDataCollectionMutex since it updates
+ /// MKernels collection.
/// \param Program a program to create kernel with.
/// \param KernelID an id of kernel to create.
/// \param KernelName a null-terminated string representing the name of kernel
@@ -151,7 +152,8 @@ class ProgramAndKernelManager {
const char *KernelName, DeviceImpl &Device);
/// Searches for kernel.
- /// This call must be protected with mutex since it reads MKernels collection.
+ /// This call must be protected with MDataCollectionMutex since it reads
+ /// MKernels collection.
/// \param KernelID an id of kernel to look for.
/// \param Device a device that kernel must be created with.
/// \return liboffload kernel for the requested configuration or nullptr if
@@ -166,10 +168,7 @@ class ProgramAndKernelManager {
std::unique_ptr<DeviceImageWrapper>>
MDeviceImageWrappers;
- // All data collections, created from data in __sycl_register_lib, must be
- // protected with this mutex. Protects data that can be modified only by
- // modules load/unload.
- std::mutex MImageCollectionMutex;
+ std::mutex MDataCollectionMutex;
// Filled by getOrCreateKernel and everything it calls inside.
std::unordered_map<
@@ -183,9 +182,6 @@ class ProgramAndKernelManager {
// Controls lifetime of programs.
std::unordered_map<ol_program_handle_t, std::unique_ptr<ProgramWrapper>>
MProgramWrappers;
- // All data collections, used and modified by kernel submissions, must be
- // protected with this mutex.
- std::mutex MKernelCollectionsMutex;
};
} // namespace detail
>From 49ee3f34a182d171afce6e0331107578184539be Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Fri, 10 Apr 2026 05:11:55 -0700
Subject: [PATCH 8/9] fix format
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/detail/program_manager.cpp | 16 ++++++++--------
1 file changed, 8 insertions(+), 8 deletions(-)
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index f461c0bcdfec1..47755bbd69a02 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -165,8 +165,9 @@ ProgramAndKernelManager::getDeviceImage(std::string_view KernelName,
"No kernel named " + std::string(KernelName) + " was found");
}
-ol_symbol_handle_t ProgramAndKernelManager::getOrCreateKernel(const char *KernelName,
- DeviceImpl &Device) {
+ol_symbol_handle_t
+ProgramAndKernelManager::getOrCreateKernel(const char *KernelName,
+ DeviceImpl &Device) {
std::unique_lock<std::mutex> ImageGuard(MDataCollectionMutex);
auto KernelIDIt = MKernelNameToID.find(KernelName);
@@ -194,7 +195,7 @@ ol_symbol_handle_t ProgramAndKernelManager::getOrCreateKernel(const char *Kernel
ol_program_handle_t
ProgramAndKernelManager::getOrCreateProgram(DeviceImpl &Device,
- DeviceImageWrapper *DevImage) {
+ DeviceImageWrapper *DevImage) {
if (auto DevToProgramIt = MPrograms.find(DevImage);
DevToProgramIt != MPrograms.end()) {
auto ProgramIt = DevToProgramIt->second.find(Device.getOLHandle());
@@ -212,10 +213,9 @@ ProgramAndKernelManager::getOrCreateProgram(DeviceImpl &Device,
return Program;
}
-ol_symbol_handle_t ProgramAndKernelManager::createKernel(ol_program_handle_t Program,
- const kernel_id &KernelID,
- const char *KernelName,
- DeviceImpl &Device) {
+ol_symbol_handle_t ProgramAndKernelManager::createKernel(
+ ol_program_handle_t Program, const kernel_id &KernelID,
+ const char *KernelName, DeviceImpl &Device) {
assert((getKernel(KernelID, Device) == nullptr) &&
"Attempt to create kernel that already exists.");
ol_symbol_handle_t Kernel{};
@@ -227,7 +227,7 @@ ol_symbol_handle_t ProgramAndKernelManager::createKernel(ol_program_handle_t Pro
}
ol_symbol_handle_t ProgramAndKernelManager::getKernel(const kernel_id &KernelID,
- DeviceImpl &Device) {
+ DeviceImpl &Device) {
auto Range = MKernels.equal_range(KernelID);
for (auto Kernels = Range.first; Kernels != Range.second; ++Kernels) {
auto &[KernelDevice, KernelSymbol] = Kernels->second;
>From 3d40e17f3e8f2f25d2bcd1d5e25c0e547f31688a Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Mon, 13 Apr 2026 06:03:54 -0700
Subject: [PATCH 9/9] redesigned kernel creation
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/detail/device_image_wrapper.hpp | 67 ++++++--
libsycl/src/detail/device_impl.hpp | 2 +-
libsycl/src/detail/device_kernel_info.hpp | 59 ++++++++
libsycl/src/detail/kernel_id.hpp | 83 ----------
libsycl/src/detail/program_manager.cpp | 160 +++++---------------
libsycl/src/detail/program_manager.hpp | 106 ++-----------
6 files changed, 172 insertions(+), 305 deletions(-)
create mode 100644 libsycl/src/detail/device_kernel_info.hpp
delete mode 100644 libsycl/src/detail/kernel_id.hpp
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
index 45dbb84a5a14a..11b0d67c4471e 100644
--- a/libsycl/src/detail/device_image_wrapper.hpp
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -19,22 +19,57 @@
#include <detail/device_binary_structures.hpp>
+#include <OffloadAPI.h>
+
+#include <map>
+
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
-/// A wrapper of __sycl_tgt_device_image structure to help with its fields
-/// parsing, iteration over data and data transformation.
-class DeviceImageWrapper {
+class DeviceImageManager;
+
+/// A wrapper of liboffload program handle to manage its lifetime.
+class ProgramWrapper {
+public:
+ /// Constructs ProgramWrapper by creating liboffload program with the provided
+ /// arguments.
+ ///
+ /// \param Device is a device to use for program creation.
+ /// \param DevImage is a device image (wrapped __sycl_tgt_device_image) to use
+ /// for program creation.
+ /// \throw sycl::exception with sycl::errc::runtime when failed to create
+ /// program.
+ ProgramWrapper(ol_device_handle_t Device, DeviceImageManager &DevImage);
+
+ /// Releases the corresponding liboffload program handle by calling
+ /// olDestroyProgram.
+ ~ProgramWrapper();
+
+ ProgramWrapper(const ProgramWrapper &) = delete;
+ ProgramWrapper &operator=(const ProgramWrapper &) = delete;
+ ProgramWrapper(ProgramWrapper &&) = delete;
+ ProgramWrapper &operator=(ProgramWrapper &&) = delete;
+
+ /// \return the corresponding liboffload program handle.
+ ol_program_handle_t getHandle() { return MProgram; }
+
+private:
+ ol_program_handle_t MProgram{};
+};
+
+/// This class manages all work with device images: from data parsing to program
+/// creation.
+class DeviceImageManager {
public:
- DeviceImageWrapper(const __sycl_tgt_device_image &Bin) : MBin(&Bin) {}
+ DeviceImageManager(const __sycl_tgt_device_image &Bin) : MBin(&Bin) {}
// Explicitly delete copy constructor/operator= to avoid unintentional copies.
- DeviceImageWrapper(const DeviceImageWrapper &) = delete;
- DeviceImageWrapper &operator=(const DeviceImageWrapper &) = delete;
+ DeviceImageManager(const DeviceImageManager &) = delete;
+ DeviceImageManager &operator=(const DeviceImageManager &) = delete;
- DeviceImageWrapper(DeviceImageWrapper &&) = default;
- DeviceImageWrapper &operator=(DeviceImageWrapper &&) = default;
+ DeviceImageManager(DeviceImageManager &&) = default;
+ DeviceImageManager &operator=(DeviceImageManager &&) = default;
- ~DeviceImageWrapper() = default;
+ ~DeviceImageManager() = default;
/// \return a reference to the corresponding raw __sycl_tgt_device_image
/// object.
@@ -45,7 +80,21 @@ class DeviceImageWrapper {
return static_cast<size_t>(MBin->ImageEnd - MBin->ImageStart);
}
+ ol_program_handle_t getOrCreateProgram(ol_device_handle_t DeviceHandle) {
+ auto ProgramIt = MPrograms.find(DeviceHandle);
+ if (ProgramIt == MPrograms.end()) {
+ ProgramIt =
+ MPrograms.emplace_hint(ProgramIt, std::piecewise_construct,
+ std::forward_as_tuple(DeviceHandle),
+ std::forward_as_tuple(DeviceHandle, *this));
+ }
+
+ return ProgramIt->second.getHandle();
+ }
+
protected:
+ std::unordered_map<ol_device_handle_t, ProgramWrapper> MPrograms;
+
const __sycl_tgt_device_image *get() const { return MBin; }
__sycl_tgt_device_image const *MBin{};
diff --git a/libsycl/src/detail/device_impl.hpp b/libsycl/src/detail/device_impl.hpp
index b7c127390afeb..345047b57f891 100644
--- a/libsycl/src/detail/device_impl.hpp
+++ b/libsycl/src/detail/device_impl.hpp
@@ -121,7 +121,7 @@ class DeviceImpl {
static_assert(false && "Info descriptor is not properly supported");
}
- ol_device_handle_t getOLHandle() { return MOffloadDevice; }
+ ol_device_handle_t getOLHandle() const { return MOffloadDevice; }
private:
ol_device_handle_t MOffloadDevice = {};
diff --git a/libsycl/src/detail/device_kernel_info.hpp b/libsycl/src/detail/device_kernel_info.hpp
new file mode 100644
index 0000000000000..89d33431a95f4
--- /dev/null
+++ b/libsycl/src/detail/device_kernel_info.hpp
@@ -0,0 +1,59 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+/// to add
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_DEVICE_KERNEL_INFO
+#define _LIBSYCL_DEVICE_KERNEL_INFO
+
+#include <sycl/__impl/detail/config.hpp>
+
+#include <OffloadAPI.h>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+namespace detail {
+
+// This class aggregates information specific to device kernels (i.e.
+// information that is uniform between different submissions of the same
+// kernel). Pointers to instances of this class are stored in header function
+// templates as a static variable to avoid repeated runtime lookup overhead.
+class DeviceKernelInfo {
+public:
+ DeviceKernelInfo(std::string_view KernelName, DeviceImageManager &DeviceImage)
+ : MName(KernelName), MDeviceImage(DeviceImage) {}
+
+ ol_symbol_handle_t getKernel(ol_device_handle_t Device) {
+ if (auto KernelIt = MBuiltKernels.find(Device);
+ KernelIt != MBuiltKernels.end())
+ return KernelIt->second;
+ return nullptr;
+ }
+
+ DeviceImageManager &getDeviceImage() { return MDeviceImage; }
+ std::string_view getName() { return MName; }
+
+ void addKernel(ol_device_handle_t Device, ol_symbol_handle_t Kernel) {
+ assert(MBuiltKernels.find(Device) != MBuiltKernels.end());
+ MBuiltKernels.insert({Device, Kernel});
+ }
+
+private:
+ std::unordered_map<ol_device_handle_t, ol_symbol_handle_t> MBuiltKernels;
+
+ std::string_view MName;
+ DeviceImageManager &MDeviceImage;
+};
+
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL_DEVICE_KERNEL_INFO
diff --git a/libsycl/src/detail/kernel_id.hpp b/libsycl/src/detail/kernel_id.hpp
deleted file mode 100644
index 421e4c6fe17d6..0000000000000
--- a/libsycl/src/detail/kernel_id.hpp
+++ /dev/null
@@ -1,83 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// 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
-/// This file contains the declaration of sycl::kernel_id and its implementation
-/// counterpart, which represent a kernel identificator.
-///
-//===----------------------------------------------------------------------===//
-
-#ifndef _LIBSYCL_KERNEL_ID
-#define _LIBSYCL_KERNEL_ID
-
-#include <sycl/__impl/detail/config.hpp>
-#include <sycl/__impl/detail/obj_utils.hpp>
-
-#include <memory>
-#include <string>
-
-_LIBSYCL_BEGIN_NAMESPACE_SYCL
-
-namespace detail {
-/// The class is the implementation counterpart for sycl::kernel_id, which
-/// represents a kernel identificator.
-class KernelIdImpl {
-public:
- KernelIdImpl(std::string_view Name) : MName(Name) {}
- KernelIdImpl() {}
- /// \return a null-terminated string representing the name of the kernel this
- /// id stands for.
- const char *get_name() { return MName.data(); }
-
-private:
- std::string MName;
-};
-} // namespace detail
-
-// TODO: It is not exported now, but is a part of SYCL spec.
-/// Kernel identifier.
-class kernel_id {
-public:
- kernel_id() = delete;
-
- kernel_id(const kernel_id &rhs) = default;
-
- kernel_id(kernel_id &&rhs) = default;
-
- kernel_id &operator=(const kernel_id &rhs) = default;
-
- kernel_id &operator=(kernel_id &&rhs) = default;
-
- friend bool operator==(const kernel_id &lhs, const kernel_id &rhs) {
- return lhs.impl == rhs.impl;
- }
-
- friend bool operator!=(const kernel_id &lhs, const kernel_id &rhs) {
- return !(lhs == rhs);
- }
-
- /// \returns a null-terminated string that contains the kernel name.
- const char *get_name() const noexcept { return impl->get_name(); }
-
-private:
- kernel_id(const char *Name);
-
- kernel_id(const std::shared_ptr<detail::KernelIdImpl> &Impl)
- : impl(std::move(Impl)) {}
-
- std::shared_ptr<detail::KernelIdImpl> impl;
- friend sycl::detail::ImplUtils;
-};
-
-_LIBSYCL_END_NAMESPACE_SYCL
-
-template <>
-struct std::hash<sycl::kernel_id>
- : public sycl::detail::HashBase<sycl::kernel_id> {};
-
-#endif // _LIBSYCL_KERNEL_ID
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index 47755bbd69a02..19e9017cc3a09 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -19,7 +19,7 @@ _LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
ProgramWrapper::ProgramWrapper(ol_device_handle_t Device,
- DeviceImageWrapper &DevImage) {
+ DeviceImageManager &DevImage) {
assert(Device);
callAndThrow(olCreateProgram, Device, DevImage.getRawData().ImageStart,
@@ -65,26 +65,22 @@ void ProgramAndKernelManager::registerFatBin(__sycl_tgt_bin_desc *FatbinDesc) {
if (EntriesB == EntriesE)
continue;
- std::unique_ptr<DeviceImageWrapper> NewImageWrapper =
- std::make_unique<DeviceImageWrapper>(RawDeviceImage);
+ std::unique_ptr<DeviceImageManager> NewImageWrapper =
+ std::make_unique<DeviceImageManager>(RawDeviceImage);
for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; ++EntriesIt) {
auto Name = EntriesIt->SymbolName;
- auto KernelIDIt = MKernelNameToID.find(Name);
- if (KernelIDIt == MKernelNameToID.end()) {
- sycl::kernel_id KernelID =
- detail::createSyclObjFromImpl<sycl::kernel_id>(
- std::make_shared<detail::KernelIdImpl>(Name));
- KernelIDIt = MKernelNameToID.insert(
- MKernelNameToID.end(),
- std::make_pair(std::string_view(Name), KernelID));
- }
- MKernelIDToDevImageJIT.insert(
- std::make_pair(KernelIDIt->second, NewImageWrapper.get()));
+ auto It = MDeviceKernelInfoMap.find(std::string_view(Name));
+ if (It == MDeviceKernelInfoMap.end()) {
+
+ std::ignore = MDeviceKernelInfoMap.emplace_hint(
+ It, std::piecewise_construct, std::forward_as_tuple(Name),
+ std::forward_as_tuple(Name, *NewImageWrapper));
+ }
}
- MDeviceImageWrappers.insert(
+ MDeviceImageManagers.insert(
std::make_pair(&RawDeviceImage, std::move(NewImageWrapper)));
}
}
@@ -100,8 +96,8 @@ void ProgramAndKernelManager::unregisterFatBin(
for (uint16_t I = 0; I < FatbinDesc->NumDeviceBinaries; ++I) {
const auto &RawDeviceImage = FatbinDesc->DeviceImages[I];
- auto DevImageIt = MDeviceImageWrappers.find(&RawDeviceImage);
- if (DevImageIt == MDeviceImageWrappers.end())
+ auto DevImageIt = MDeviceImageManagers.find(&RawDeviceImage);
+ if (DevImageIt == MDeviceImageManagers.end())
continue;
const llvm::offloading::EntryTy *EntriesB = RawDeviceImage.EntriesBegin;
@@ -111,134 +107,60 @@ void ProgramAndKernelManager::unregisterFatBin(
continue;
for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; ++EntriesIt) {
- if (auto KernelIDIt = MKernelNameToID.find(EntriesIt->SymbolName);
- KernelIDIt != MKernelNameToID.end()) {
- MKernelIDToDevImageJIT.erase(KernelIDIt->second);
- MKernels.erase(KernelIDIt->second);
- MKernelNameToID.erase(KernelIDIt);
+ if (auto KernelIt = MDeviceKernelInfoMap.find(EntriesIt->SymbolName);
+ KernelIt != MDeviceKernelInfoMap.end()) {
+ // Programs are attached to image and will be release with image
+ // destruction. Clear only kernel specific data by destroying its kernel
+ // info object.
+ MDeviceKernelInfoMap.erase(KernelIt);
}
}
- if (auto ProgramIt = MPrograms.find(DevImageIt->second.get());
- ProgramIt != MPrograms.end()) {
- for (auto &[Device, Program] : ProgramIt->second) {
- MProgramWrappers.erase(Program);
- }
- MPrograms.erase(ProgramIt);
- }
-
- MDeviceImageWrappers.erase(DevImageIt);
+ MDeviceImageManagers.erase(DevImageIt);
}
}
-static bool isImageTargetCompatible(const DeviceImageWrapper &Image,
- const DeviceImpl &Device) {
+static bool isImageCompatible(const DeviceImageManager &Image,
+ const DeviceImpl &Device) {
sycl::backend BE = Device.getBackend();
const char *Target = Image.getRawData().TripleString;
- return (strcmp(Target, DeviceBinaryTripleSPIRV64) == 0) &&
- (BE == sycl::backend::level_zero);
-}
+ if (!(strcmp(Target, DeviceBinaryTripleSPIRV64) == 0) &&
+ (BE == sycl::backend::level_zero))
+ return false;
-DeviceImageWrapper *
-ProgramAndKernelManager::getDeviceImage(std::string_view KernelName,
- const kernel_id &KernelID,
- DeviceImpl &Device) {
- std::lock_guard<std::mutex> Guard(MDataCollectionMutex);
- auto [Begin, End] = MKernelIDToDevImageJIT.equal_range(KernelID);
- if (Begin != End) {
- bool IsValid{};
- // TODO: with AOT (not implemented yet), we need to analyze and check
- // olIsValidBinary for AOT binaries first.
- for (auto It = Begin; It != End; ++It) {
- if (isImageTargetCompatible(*It->second, Device)) {
- callAndThrow(olIsValidBinary, Device.getOLHandle(),
- It->second->getRawData().ImageStart, It->second->getSize(),
- &IsValid);
- if (IsValid)
- return It->second;
- }
- }
- }
-
- throw exception(make_error_code(errc::runtime),
- "No kernel named " + std::string(KernelName) + " was found");
+ bool IsValid{};
+ callAndThrow(olIsValidBinary, Device.getOLHandle(),
+ Image.getRawData().ImageStart, Image.getSize(), &IsValid);
+ return IsValid;
}
ol_symbol_handle_t
-ProgramAndKernelManager::getOrCreateKernel(const char *KernelName,
+ProgramAndKernelManager::getOrCreateKernel(DeviceKernelInfo &KernelInfo,
DeviceImpl &Device) {
- std::unique_lock<std::mutex> ImageGuard(MDataCollectionMutex);
-
- auto KernelIDIt = MKernelNameToID.find(KernelName);
- if (KernelIDIt == MKernelNameToID.end())
- throw exception(make_error_code(errc::runtime),
- "No kernel named " + std::string(KernelName) +
- " was found");
std::lock_guard<std::mutex> KernelGuard(MDataCollectionMutex);
- auto Kernel = getKernel(KernelIDIt->second, Device);
- if (Kernel)
+ if (auto Kernel = KernelInfo.getKernel(Device.getOLHandle()))
return Kernel;
- DeviceImageWrapper *DevImage =
- getDeviceImage(KernelName, KernelIDIt->second, Device);
-
- ol_program_handle_t Program = getOrCreateProgram(Device, DevImage);
- assert(Program);
-
- Kernel = createKernel(Program, KernelIDIt->second, KernelName, Device);
- assert(Kernel);
- return Kernel;
-}
-
-ol_program_handle_t
-ProgramAndKernelManager::getOrCreateProgram(DeviceImpl &Device,
- DeviceImageWrapper *DevImage) {
- if (auto DevToProgramIt = MPrograms.find(DevImage);
- DevToProgramIt != MPrograms.end()) {
- auto ProgramIt = DevToProgramIt->second.find(Device.getOLHandle());
- if (ProgramIt != DevToProgramIt->second.end())
- return ProgramIt->second;
- }
+ auto &DeviceImage = KernelInfo.getDeviceImage();
- std::unique_ptr<ProgramWrapper> NewProgramWrapper(
- new ProgramWrapper(Device.getOLHandle(), *DevImage));
- auto Program = NewProgramWrapper->getHandle();
- MPrograms[DevImage].insert(std::make_pair(Device.getOLHandle(), Program));
- MProgramWrappers.insert(std::make_pair(NewProgramWrapper->getHandle(),
- std::move(NewProgramWrapper)));
+ if (!isImageCompatible(DeviceImage, Device))
+ throw exception(make_error_code(errc::runtime),
+ std::string("No compatible image for ") +
+ KernelInfo.getName().data() + " was found");
- return Program;
-}
+ auto DeviceHandle = Device.getOLHandle();
+ auto Program = DeviceImage.getOrCreateProgram(DeviceHandle);
-ol_symbol_handle_t ProgramAndKernelManager::createKernel(
- ol_program_handle_t Program, const kernel_id &KernelID,
- const char *KernelName, DeviceImpl &Device) {
- assert((getKernel(KernelID, Device) == nullptr) &&
- "Attempt to create kernel that already exists.");
ol_symbol_handle_t Kernel{};
- callAndThrow(olGetSymbol, Program, KernelName, OL_SYMBOL_KIND_KERNEL,
- &Kernel);
- MKernels.insert(
- std::make_pair(KernelID, std::make_pair(Device.getOLHandle(), Kernel)));
+ callAndThrow(olGetSymbol, Program, KernelInfo.getName().data(),
+ OL_SYMBOL_KIND_KERNEL, &Kernel);
+ KernelInfo.addKernel(DeviceHandle, Kernel);
return Kernel;
}
-ol_symbol_handle_t ProgramAndKernelManager::getKernel(const kernel_id &KernelID,
- DeviceImpl &Device) {
- auto Range = MKernels.equal_range(KernelID);
- for (auto Kernels = Range.first; Kernels != Range.second; ++Kernels) {
- auto &[KernelDevice, KernelSymbol] = Kernels->second;
- if (KernelDevice == Device.getOLHandle()) {
- assert(KernelSymbol && "Built kernel symbol can't be null");
- return KernelSymbol;
- }
- }
- return nullptr;
-}
-
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index 5ab828b25a470..7cee3b999531d 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -19,7 +19,7 @@
#include <detail/device_binary_structures.hpp>
#include <detail/device_image_wrapper.hpp>
-#include <detail/kernel_id.hpp>
+#include <detail/device_kernel_info.hpp>
#include <OffloadAPI.h>
@@ -47,35 +47,6 @@ namespace detail {
class DeviceImpl;
-/// A wrapper of liboffload program handle to manage its lifetime.
-class ProgramWrapper {
-public:
- /// Constructs ProgramWrapper by creating liboffload program with the provided
- /// arguments.
- ///
- /// \param Device is a device to use for program creation.
- /// \param DevImage is a device image (wrapped __sycl_tgt_device_image) to use
- /// for program creation.
- /// \throw sycl::exception with sycl::errc::runtime when failed to create
- /// program.
- ProgramWrapper(ol_device_handle_t Device, DeviceImageWrapper &DevImage);
-
- /// Releases the corresponding liboffload program handle by calling
- /// olDestroyProgram.
- ~ProgramWrapper();
-
- ProgramWrapper(const ProgramWrapper &) = delete;
- ProgramWrapper &operator=(const ProgramWrapper &) = delete;
- ProgramWrapper(ProgramWrapper &&) = delete;
- ProgramWrapper &operator=(ProgramWrapper &&) = delete;
-
- /// \return the corresponding liboffload program handle.
- ol_program_handle_t getHandle() { return MProgram; }
-
-private:
- ol_program_handle_t MProgram{};
-};
-
/// A class to manage programs and kernels.
class ProgramAndKernelManager {
@@ -102,12 +73,12 @@ class ProgramAndKernelManager {
/// Creates liboffload kernel that is ready for execution.
/// Thread-safe.
- /// \param KernelName a null-terminated string representing a name of kernel
- /// to be created.
+ /// \param KernelInfo a set of kernel specific data: name, corresponding
+ /// device image, etc.
/// \param Device a device for which this kernel must be compiled.
/// \return liboffload kernel handle that is ready to be passed to kernel
/// execution methods.
- ol_symbol_handle_t getOrCreateKernel(const char *KernelName,
+ ol_symbol_handle_t getOrCreateKernel(DeviceKernelInfo &KernelInfo,
DeviceImpl &Device);
private:
@@ -116,72 +87,21 @@ class ProgramAndKernelManager {
ProgramAndKernelManager(ProgramAndKernelManager const &) = delete;
ProgramAndKernelManager &operator=(ProgramAndKernelManager const &) = delete;
- /// Searches for a device image that contains the requested kernel and is
- /// compatible with the requested device.
- /// This call must be protected with MDataCollectionMutex.
- /// \param KernelName a null-terminated string representing the name of the
- /// kernel to obtain a device image for.
- /// \param KernelID a kernel id matching KernelName.
- /// \param Device a device with which device image must be compatible.
- /// \throw sycl::exception with sycl::errc::runtime if the device image
- /// validation failed in liboffload or if no compatible image was found.
- DeviceImageWrapper *getDeviceImage(std::string_view KernelName,
- const kernel_id &KernelID,
- DeviceImpl &Device);
-
- /// Searches for or creates a program.
- /// This call must be protected with MDataCollectionMutex since it updates
- /// MPrograms and MProgramWrappers collections.
- /// \param Device a device that program must be created with.
- /// \param DevImage a device image to get or create program with.
- /// \return liboffload program for the requested configuration.
- ol_program_handle_t getOrCreateProgram(DeviceImpl &Device,
- DeviceImageWrapper *DevImage);
-
- /// Creates kernel from program.
- /// This call must be protected with MDataCollectionMutex since it updates
- /// MKernels collection.
- /// \param Program a program to create kernel with.
- /// \param KernelID an id of kernel to create.
- /// \param KernelName a null-terminated string representing the name of kernel
- /// to create.
- /// \param Device a device that kernel must be created with.
- /// \return liboffload kernel for the requested configuration.
- ol_symbol_handle_t createKernel(ol_program_handle_t Program,
- const kernel_id &KernelID,
- const char *KernelName, DeviceImpl &Device);
-
- /// Searches for kernel.
- /// This call must be protected with MDataCollectionMutex since it reads
- /// MKernels collection.
- /// \param KernelID an id of kernel to look for.
- /// \param Device a device that kernel must be created with.
- /// \return liboffload kernel for the requested configuration or nullptr if
- /// such kernel is not found.
- ol_symbol_handle_t getKernel(const kernel_id &KernelID, DeviceImpl &Device);
-
// Filled by registerFatBin(...).
- std::unordered_map<std::string_view, kernel_id> MKernelNameToID;
- std::unordered_map<kernel_id, DeviceImageWrapper *> MKernelIDToDevImageJIT;
- // Controls lifetime of device image ptr and wrapper.
- std::unordered_map<const __sycl_tgt_device_image *,
- std::unique_ptr<DeviceImageWrapper>>
- MDeviceImageWrappers;
-
- std::mutex MDataCollectionMutex;
+ // Map for storing device kernel information. Runtime lookup should be avoided
+ // by caching the pointers when possible.
+ std::unordered_map<std::string_view, DeviceKernelInfo> MDeviceKernelInfoMap;
- // Filled by getOrCreateKernel and everything it calls inside.
- std::unordered_map<
- DeviceImageWrapper *,
- std::unordered_map<ol_device_handle_t, ol_program_handle_t>>
- MPrograms;
- std::unordered_multimap<kernel_id,
- std::pair<ol_device_handle_t, ol_symbol_handle_t>>
- MKernels;
+ // Controls lifetime of device images.
+ std::unordered_map<const __sycl_tgt_device_image *,
+ std::unique_ptr<DeviceImageManager>>
+ MDeviceImageManagers;
// Controls lifetime of programs.
std::unordered_map<ol_program_handle_t, std::unique_ptr<ProgramWrapper>>
MProgramWrappers;
+
+ std::mutex MDataCollectionMutex;
};
} // namespace detail
More information about the llvm-commits
mailing list