[llvm] [libsycl] add single_task (PR #192499)
Kseniya Tikhomirova via llvm-commits
llvm-commits at lists.llvm.org
Fri Apr 17 04:28:12 PDT 2026
https://github.com/KseniyaTikhomirova updated https://github.com/llvm/llvm-project/pull/192499
>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 01/13] [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 02/13] 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 03/13] 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 04/13] [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 05/13] [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 06/13] 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 07/13] 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 08/13] 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 09/13] 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
>From 32c3481ddf8956d0ee43f430b93ad352b9821412 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Tue, 14 Apr 2026 02:56:59 -0700
Subject: [PATCH 10/13] code cleanup
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/detail/device_image_wrapper.hpp | 13 ++++--
libsycl/src/detail/device_impl.hpp | 1 +
libsycl/src/detail/device_kernel_info.hpp | 49 +++++++++++++++------
libsycl/src/detail/program_manager.cpp | 2 +-
libsycl/src/detail/program_manager.hpp | 6 +--
5 files changed, 48 insertions(+), 23 deletions(-)
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
index 11b0d67c4471e..d2b952acb39ed 100644
--- a/libsycl/src/detail/device_image_wrapper.hpp
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -7,8 +7,8 @@
//===----------------------------------------------------------------------===//
///
/// \file
-/// This file contains the declaration of the helper for raw device image
-/// parsing and iteration.
+/// This file contains the declaration of the helpers for device images and
+/// programs.
///
//===----------------------------------------------------------------------===//
@@ -51,7 +51,7 @@ class ProgramWrapper {
ProgramWrapper &operator=(ProgramWrapper &&) = delete;
/// \return the corresponding liboffload program handle.
- ol_program_handle_t getHandle() { return MProgram; }
+ ol_program_handle_t getOLHandle() { return MProgram; }
private:
ol_program_handle_t MProgram{};
@@ -80,6 +80,11 @@ class DeviceImageManager {
return static_cast<size_t>(MBin->ImageEnd - MBin->ImageStart);
}
+ /// Returns liboffload program handle by lookup of existing programs or by
+ /// creation of a new one from this image.
+ /// \param DeviceHandle liboffload handle of device the program must be
+ /// compatible with.
+ /// \return liboffload handle of the program compatible with specified device.
ol_program_handle_t getOrCreateProgram(ol_device_handle_t DeviceHandle) {
auto ProgramIt = MPrograms.find(DeviceHandle);
if (ProgramIt == MPrograms.end()) {
@@ -89,7 +94,7 @@ class DeviceImageManager {
std::forward_as_tuple(DeviceHandle, *this));
}
- return ProgramIt->second.getHandle();
+ return ProgramIt->second.getOLHandle();
}
protected:
diff --git a/libsycl/src/detail/device_impl.hpp b/libsycl/src/detail/device_impl.hpp
index 345047b57f891..f5012fe84c069 100644
--- a/libsycl/src/detail/device_impl.hpp
+++ b/libsycl/src/detail/device_impl.hpp
@@ -121,6 +121,7 @@ class DeviceImpl {
static_assert(false && "Info descriptor is not properly supported");
}
+ /// \return the corresponding liboffload device handle.
ol_device_handle_t getOLHandle() const { return MOffloadDevice; }
private:
diff --git a/libsycl/src/detail/device_kernel_info.hpp b/libsycl/src/detail/device_kernel_info.hpp
index 89d33431a95f4..6c9928f711ee2 100644
--- a/libsycl/src/detail/device_kernel_info.hpp
+++ b/libsycl/src/detail/device_kernel_info.hpp
@@ -7,7 +7,9 @@
//===----------------------------------------------------------------------===//
///
/// \file
-/// to add
+/// This file contains the declaration of the class that aggregates information
+/// specific to device kernels (i.e. information that is uniform between
+/// different submissions of the same kernel).
///
//===----------------------------------------------------------------------===//
@@ -21,35 +23,54 @@
_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 ProgramAndKernelManager;
+
+// 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:
+ /// Constructs device kernel info instance.
+ ///
+ /// \param KernelName a name of kernel.
+ /// \param DeviceImage a device image containing device code of this kernel.
DeviceKernelInfo(std::string_view KernelName, DeviceImageManager &DeviceImage)
: MName(KernelName), MDeviceImage(DeviceImage) {}
- ol_symbol_handle_t getKernel(ol_device_handle_t Device) {
+ /// \return the name of this kernel.
+ std::string_view getName() { return MName; }
+
+private:
+ std::unordered_map<ol_device_handle_t, ol_symbol_handle_t> MBuiltKernels;
+
+ std::string_view MName;
+ DeviceImageManager &MDeviceImage;
+
+ /// Searches for the existing kernel handle compatible with the specified
+ /// device.
+ /// \param Device a device the kernel must be compatible with.
+ /// \return a liboffload kernel handle if and only if built kernel was found,
+ /// otherwise returns nullptr.
+ ol_symbol_handle_t getKernel(ol_device_handle_t Device) const {
if (auto KernelIt = MBuiltKernels.find(Device);
KernelIt != MBuiltKernels.end())
return KernelIt->second;
return nullptr;
}
- DeviceImageManager &getDeviceImage() { return MDeviceImage; }
- std::string_view getName() { return MName; }
+ /// \return device image which contains device code of this kernel.
+ DeviceImageManager &getDeviceImage() const { return MDeviceImage; }
+ /// Attaches liboffload kernel handle to this device kernel info object.
+ /// \param Device the device the kernel symbol was created for.
+ /// \param Kernel the liboffload kernel symbol to attach..
void addKernel(ol_device_handle_t Device, ol_symbol_handle_t Kernel) {
- assert(MBuiltKernels.find(Device) != MBuiltKernels.end());
+ assert(Kernel && Device &&
+ 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;
+ /// Kernel info update is intended to be done only by ProgramAndKernelManager.
+ friend class ProgramAndKernelManager;
};
} // namespace detail
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index 19e9017cc3a09..7cec508921279 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -109,7 +109,7 @@ void ProgramAndKernelManager::unregisterFatBin(
for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; ++EntriesIt) {
if (auto KernelIt = MDeviceKernelInfoMap.find(EntriesIt->SymbolName);
KernelIt != MDeviceKernelInfoMap.end()) {
- // Programs are attached to image and will be release with image
+ // Programs are attached to image and will be released with image
// destruction. Clear only kernel specific data by destroying its kernel
// info object.
MDeviceKernelInfoMap.erase(KernelIt);
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index 7cee3b999531d..d3b64c154fd2f 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -97,10 +97,8 @@ class ProgramAndKernelManager {
std::unique_ptr<DeviceImageManager>>
MDeviceImageManagers;
- // Controls lifetime of programs.
- std::unordered_map<ol_program_handle_t, std::unique_ptr<ProgramWrapper>>
- MProgramWrappers;
-
+ // All work with device images and data related to it must be wrapped with a
+ // lock of this mutex.
std::mutex MDataCollectionMutex;
};
>From 22cf0f4a8725f66c076cd20a7301f11d1af72fef Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 15 Apr 2026 09:34:11 -0700
Subject: [PATCH 11/13] fix comments
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/CMakeLists.txt | 1 +
libsycl/src/detail/device_image_wrapper.cpp | 39 +++++++++++++++++++++
libsycl/src/detail/device_image_wrapper.hpp | 37 ++++++++-----------
libsycl/src/detail/device_kernel_info.hpp | 23 ++++++------
libsycl/src/detail/program_manager.cpp | 26 ++++----------
libsycl/src/detail/program_manager.hpp | 6 ++--
6 files changed, 76 insertions(+), 56 deletions(-)
create mode 100644 libsycl/src/detail/device_image_wrapper.cpp
diff --git a/libsycl/src/CMakeLists.txt b/libsycl/src/CMakeLists.txt
index 7b9826fb8a3de..5b75db713353a 100644
--- a/libsycl/src/CMakeLists.txt
+++ b/libsycl/src/CMakeLists.txt
@@ -95,6 +95,7 @@ set(LIBSYCL_SOURCES
"usm_functions.cpp"
"detail/context_impl.cpp"
"detail/event_impl.cpp"
+ "detail/device_image_wrapper.cpp"
"detail/device_impl.cpp"
"detail/global_objects.cpp"
"detail/platform_impl.cpp"
diff --git a/libsycl/src/detail/device_image_wrapper.cpp b/libsycl/src/detail/device_image_wrapper.cpp
new file mode 100644
index 0000000000000..a150feb6516e6
--- /dev/null
+++ b/libsycl/src/detail/device_image_wrapper.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/device_image_wrapper.hpp>
+
+#include <detail/offload/offload_utils.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+namespace detail {
+
+ProgramWrapper::ProgramWrapper(ol_device_handle_t Device,
+ DeviceImageManager &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.
+}
+
+ol_program_handle_t
+DeviceImageManager::getOrCreateProgram(ol_device_handle_t DeviceHandle) {
+ const auto &[Iterator, Flag] = MPrograms.emplace(
+ std::piecewise_construct, std::forward_as_tuple(DeviceHandle),
+ std::forward_as_tuple(DeviceHandle, *this));
+ return Iterator->second.getOLHandle();
+}
+
+} // namespace detail
+_LIBSYCL_END_NAMESPACE_SYCL
\ No newline at end of file
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
index d2b952acb39ed..2695264f89d2b 100644
--- a/libsycl/src/detail/device_image_wrapper.hpp
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -21,7 +21,7 @@
#include <OffloadAPI.h>
-#include <map>
+#include <unordered_map>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
@@ -31,13 +31,13 @@ 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.
+ /// Constructs ProgramWrapper by creating a 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
+ /// \param Device is the device to use for program creation.
+ /// \param DevImage is the device image (wrapped __sycl_tgt_device_image) to
+ /// use for program creation.
+ /// \throw sycl::exception with sycl::errc::runtime when failed to create the
/// program.
ProgramWrapper(ol_device_handle_t Device, DeviceImageManager &DevImage);
@@ -80,22 +80,13 @@ class DeviceImageManager {
return static_cast<size_t>(MBin->ImageEnd - MBin->ImageStart);
}
- /// Returns liboffload program handle by lookup of existing programs or by
- /// creation of a new one from this image.
- /// \param DeviceHandle liboffload handle of device the program must be
- /// compatible with.
- /// \return liboffload handle of the program compatible with specified device.
- 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.getOLHandle();
- }
+ /// Returns a liboffload program handle by looking up existing programs or
+ /// creating a new one from this image.
+ /// \param DeviceHandle the liboffload handle of the device the program must
+ /// be compatible with.
+ /// \return the liboffload handle of the program compatible with the specified
+ /// device.
+ ol_program_handle_t getOrCreateProgram(ol_device_handle_t DeviceHandle);
protected:
std::unordered_map<ol_device_handle_t, ProgramWrapper> MPrograms;
diff --git a/libsycl/src/detail/device_kernel_info.hpp b/libsycl/src/detail/device_kernel_info.hpp
index 6c9928f711ee2..2d5fdf9fe9ee9 100644
--- a/libsycl/src/detail/device_kernel_info.hpp
+++ b/libsycl/src/detail/device_kernel_info.hpp
@@ -25,14 +25,15 @@ namespace detail {
class ProgramAndKernelManager;
-// Pointers to instances of this class are stored in header function templates
-// as a static variable to avoid repeated runtime lookup overhead.
+// TODO: Pointers to instances of this class are supported to be stored in
+// header function templates as a static variable to avoid repeated runtime
+// lookup overhead.
class DeviceKernelInfo {
public:
- /// Constructs device kernel info instance.
+ /// Constructs a device kernel info instance.
///
- /// \param KernelName a name of kernel.
- /// \param DeviceImage a device image containing device code of this kernel.
+ /// \param KernelName the name of the kernel.
+ /// \param DeviceImage the device image containing device code of this kernel.
DeviceKernelInfo(std::string_view KernelName, DeviceImageManager &DeviceImage)
: MName(KernelName), MDeviceImage(DeviceImage) {}
@@ -47,9 +48,9 @@ class DeviceKernelInfo {
/// Searches for the existing kernel handle compatible with the specified
/// device.
- /// \param Device a device the kernel must be compatible with.
- /// \return a liboffload kernel handle if and only if built kernel was found,
- /// otherwise returns nullptr.
+ /// \param Device the device the kernel must be compatible with.
+ /// \return a liboffload kernel handle if a built kernel was found; otherwise
+ /// returns nullptr.
ol_symbol_handle_t getKernel(ol_device_handle_t Device) const {
if (auto KernelIt = MBuiltKernels.find(Device);
KernelIt != MBuiltKernels.end())
@@ -57,12 +58,12 @@ class DeviceKernelInfo {
return nullptr;
}
- /// \return device image which contains device code of this kernel.
+ /// \return the device image containing the device code of this kernel.
DeviceImageManager &getDeviceImage() const { return MDeviceImage; }
- /// Attaches liboffload kernel handle to this device kernel info object.
+ /// Attaches a liboffload kernel handle to this device kernel info object.
/// \param Device the device the kernel symbol was created for.
- /// \param Kernel the liboffload kernel symbol to attach..
+ /// \param Kernel the liboffload kernel symbol to attach.
void addKernel(ol_device_handle_t Device, ol_symbol_handle_t Kernel) {
assert(Kernel && Device &&
MBuiltKernels.find(Device) == MBuiltKernels.end());
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index 7cec508921279..90d7c48d3d1c7 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -18,20 +18,6 @@
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
-ProgramWrapper::ProgramWrapper(ol_device_handle_t Device,
- DeviceImageManager &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;
}
@@ -74,9 +60,11 @@ void ProgramAndKernelManager::registerFatBin(__sycl_tgt_bin_desc *FatbinDesc) {
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));
+ [[maybe_unused]] auto [Iterator, EmplaceSucceeded] =
+ MDeviceKernelInfoMap.emplace(
+ std::piecewise_construct, std::forward_as_tuple(Name),
+ std::forward_as_tuple(Name, *NewImageWrapper));
+ assert(EmplaceSucceeded && "Kernel name found in multiple images");
}
}
@@ -125,8 +113,8 @@ static bool isImageCompatible(const DeviceImageManager &Image,
sycl::backend BE = Device.getBackend();
const char *Target = Image.getRawData().TripleString;
- if (!(strcmp(Target, DeviceBinaryTripleSPIRV64) == 0) &&
- (BE == sycl::backend::level_zero))
+ if (!(strcmp(Target, DeviceBinaryTripleSPIRV64) == 0 &&
+ BE == sycl::backend::level_zero))
return false;
bool IsValid{};
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index d3b64c154fd2f..60b781e3b66e5 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -71,12 +71,12 @@ class ProgramAndKernelManager {
/// data passed to registerFatBin.
void unregisterFatBin(__sycl_tgt_bin_desc *FatbinDesc);
- /// Creates liboffload kernel that is ready for execution.
+ /// Creates a liboffload kernel that is ready for execution.
/// Thread-safe.
/// \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
+ /// \param Device the device for which this kernel must be compiled.
+ /// \return a liboffload kernel handle that is ready to be passed to kernel
/// execution methods.
ol_symbol_handle_t getOrCreateKernel(DeviceKernelInfo &KernelInfo,
DeviceImpl &Device);
>From dd46e5f2dfb639300f72fc0de1dba60a36c6779e Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Thu, 16 Apr 2026 05:12:00 -0700
Subject: [PATCH 12/13] fix comments
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/src/detail/device_image_wrapper.cpp | 2 +-
libsycl/src/detail/device_image_wrapper.hpp | 5 +++--
libsycl/src/detail/device_kernel_info.hpp | 14 ++++++++------
libsycl/src/detail/program_manager.hpp | 2 +-
4 files changed, 13 insertions(+), 10 deletions(-)
diff --git a/libsycl/src/detail/device_image_wrapper.cpp b/libsycl/src/detail/device_image_wrapper.cpp
index a150feb6516e6..36a7be0e67f87 100644
--- a/libsycl/src/detail/device_image_wrapper.cpp
+++ b/libsycl/src/detail/device_image_wrapper.cpp
@@ -36,4 +36,4 @@ DeviceImageManager::getOrCreateProgram(ol_device_handle_t DeviceHandle) {
}
} // namespace detail
-_LIBSYCL_END_NAMESPACE_SYCL
\ No newline at end of file
+_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
index 2695264f89d2b..73bdd8e7f7df1 100644
--- a/libsycl/src/detail/device_image_wrapper.hpp
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -80,8 +80,9 @@ class DeviceImageManager {
return static_cast<size_t>(MBin->ImageEnd - MBin->ImageStart);
}
- /// Returns a liboffload program handle by looking up existing programs or
- /// creating a new one from this image.
+ /// Returns a liboffload program which is compatible with the specified
+ /// device. Searches among existing programs and creates a new one if no
+ /// compatible image is found.
/// \param DeviceHandle the liboffload handle of the device the program must
/// be compatible with.
/// \return the liboffload handle of the program compatible with the specified
diff --git a/libsycl/src/detail/device_kernel_info.hpp b/libsycl/src/detail/device_kernel_info.hpp
index 2d5fdf9fe9ee9..148e2e366777d 100644
--- a/libsycl/src/detail/device_kernel_info.hpp
+++ b/libsycl/src/detail/device_kernel_info.hpp
@@ -52,10 +52,10 @@ class DeviceKernelInfo {
/// \return a liboffload kernel handle if a built kernel was found; otherwise
/// returns nullptr.
ol_symbol_handle_t getKernel(ol_device_handle_t Device) const {
- if (auto KernelIt = MBuiltKernels.find(Device);
- KernelIt != MBuiltKernels.end())
- return KernelIt->second;
- return nullptr;
+ auto KernelIt = MBuiltKernels.find(Device);
+ if (KernelIt == MBuiltKernels.end())
+ return nullptr;
+ return KernelIt->second;
}
/// \return the device image containing the device code of this kernel.
@@ -65,8 +65,10 @@ class DeviceKernelInfo {
/// \param Device the device the kernel symbol was created for.
/// \param Kernel the liboffload kernel symbol to attach.
void addKernel(ol_device_handle_t Device, ol_symbol_handle_t Kernel) {
- assert(Kernel && Device &&
- MBuiltKernels.find(Device) == MBuiltKernels.end());
+ assert(Kernel && "Invalid liboffload kernel handle");
+ assert(Device && "Invalid liboffload device handle");
+ assert((MBuiltKernels.find(Device) == MBuiltKernels.end()) &&
+ "Kernel is being managed already");
MBuiltKernels.insert({Device, Kernel});
}
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index 60b781e3b66e5..f5d3a1c8dc6dd 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -72,7 +72,7 @@ class ProgramAndKernelManager {
void unregisterFatBin(__sycl_tgt_bin_desc *FatbinDesc);
/// Creates a liboffload kernel that is ready for execution.
- /// Thread-safe.
+ /// This method is thread-safe (protected with MDataCollectionMutex).
/// \param KernelInfo a set of kernel specific data: name, corresponding
/// device image, etc.
/// \param Device the device for which this kernel must be compiled.
>From 79372977bc5f059e5d7602d6ca1cccc52798b702 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 15 Apr 2026 05:55:42 -0700
Subject: [PATCH 13/13] single_task on top of getKernelInfo
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/docs/index.rst | 4 +
.../__impl/detail/get_device_kernel_info.hpp | 43 +++++++
.../sycl/__impl/detail/unified_range_view.hpp | 51 ++++++++
libsycl/include/sycl/__impl/queue.hpp | 99 ++++++++++++++++
libsycl/src/detail/global_objects.cpp | 2 +
libsycl/src/detail/program_manager.cpp | 17 +++
libsycl/src/detail/program_manager.hpp | 6 +
libsycl/src/detail/queue_impl.cpp | 111 ++++++++++++++++++
libsycl/src/detail/queue_impl.hpp | 39 ++++++
libsycl/src/queue.cpp | 19 +++
10 files changed, 391 insertions(+)
create mode 100644 libsycl/include/sycl/__impl/detail/get_device_kernel_info.hpp
create mode 100644 libsycl/include/sycl/__impl/detail/unified_range_view.hpp
diff --git a/libsycl/docs/index.rst b/libsycl/docs/index.rst
index 9aa36b4a54c57..04691a96a188a 100644
--- a/libsycl/docs/index.rst
+++ b/libsycl/docs/index.rst
@@ -113,6 +113,10 @@ TODO for added SYCL classes
* to implement submit & copy with accessors (low priority)
* get_info & properties
* ctors that accepts context (blocked by lack of liboffload support)
+ * nd_range kernel submissions
+ * cross-context events wait (host tasks are needed)
+ * implement check if lambda arguments are device copyable (requires clang support of corresponding builtins) unless FE will fully cover it
+ * kernel instantiating on host (debugging purposes)
* ``property_list``: to fully implement and integrate with existing SYCL runtime classes supporting it
* usm allocations:
diff --git a/libsycl/include/sycl/__impl/detail/get_device_kernel_info.hpp b/libsycl/include/sycl/__impl/detail/get_device_kernel_info.hpp
new file mode 100644
index 0000000000000..292755037410e
--- /dev/null
+++ b/libsycl/include/sycl/__impl/detail/get_device_kernel_info.hpp
@@ -0,0 +1,43 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 helper function to query kernel info that is uniform
+/// between different submissions of the same kernel.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_GET_DEV_INFO
+#define _LIBSYCL_GET_DEV_INFO
+
+#include <sycl/__impl/detail/config.hpp>
+
+#include <string_view>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+class DeviceKernelInfo;
+// Lifetime of the underlying `DeviceKernelInfo` is tied to the availability of
+// the `sycl_device_binaries` corresponding to this kernel. In other words, once
+// user library is unloaded (see __sycl_unregister_lib), program manager
+// destroys this `DeviceKernelInfo` object and the reference returned from here
+// becomes stale.
+_LIBSYCL_EXPORT DeviceKernelInfo &getDeviceKernelInfo(std::string_view);
+
+template <class KernelName>
+DeviceKernelInfo &getDeviceKernelInfo(std::string_view KernelNameStr) {
+ static DeviceKernelInfo &Info = getDeviceKernelInfo(KernelNameStr);
+ return Info;
+}
+
+} // namespace detail
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL_GET_DEV_INFO
diff --git a/libsycl/include/sycl/__impl/detail/unified_range_view.hpp b/libsycl/include/sycl/__impl/detail/unified_range_view.hpp
new file mode 100644
index 0000000000000..4bcaa48eec757
--- /dev/null
+++ b/libsycl/include/sycl/__impl/detail/unified_range_view.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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains helper function class to unify ABI for different kernel
+/// ranges.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_DETAIL_UNIFIED_RANGE_VIEW_HPP
+#define _LIBSYCL___IMPL_DETAIL_UNIFIED_RANGE_VIEW_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+/// The structure to keep dimension and references to ranges unified for
+/// all dimensions.
+struct UnifiedRangeView {
+ /// Default contructed view matches the single task execution range.
+ UnifiedRangeView() = default;
+ UnifiedRangeView(const UnifiedRangeView &Desc) = default;
+ UnifiedRangeView(UnifiedRangeView &&Desc) = default;
+ UnifiedRangeView &operator=(const UnifiedRangeView &Desc) = default;
+ UnifiedRangeView &operator=(UnifiedRangeView &&Desc) = default;
+ ~UnifiedRangeView() = default;
+
+ // TODO: ctors with sycl::range and nd::range will be added later.
+
+ UnifiedRangeView(const size_t *GlobalSize, const size_t *LocalSize,
+ const size_t *Offset, size_t Dims)
+ : MGlobalSize(GlobalSize), MLocalSize(LocalSize), MOffset(Offset),
+ MDims(Dims) {}
+
+ const size_t *MGlobalSize = nullptr;
+ const size_t *MLocalSize = nullptr;
+ const size_t *MOffset = nullptr;
+ size_t MDims = 1;
+};
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_DETAIL_UNIFIED_RANGE_VIEW_HPP
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 41b018b681b8e..991ebe7f2a666 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -17,11 +17,14 @@
#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>
#include <sycl/__impl/detail/default_async_handler.hpp>
+#include <sycl/__impl/detail/get_device_kernel_info.hpp>
#include <sycl/__impl/detail/obj_utils.hpp>
+#include <sycl/__impl/detail/unified_range_view.hpp>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
@@ -29,6 +32,27 @@ class context;
namespace detail {
class QueueImpl;
+
+template <typename, typename T> struct CheckFunctionSignature {
+ static_assert(std::integral_constant<T, false>::value,
+ "Second template parameter is required to be of function type");
+};
+
+template <typename F, typename RetT, typename... Args>
+struct CheckFunctionSignature<F, RetT(Args...)> {
+private:
+ template <typename T>
+ static constexpr auto check(T *) -> typename std::is_same<
+ decltype(std::declval<T>().operator()(std::declval<Args>()...)),
+ RetT>::type;
+
+ template <typename> static constexpr std::false_type check(...);
+
+ using type = decltype(check<F>(0));
+
+public:
+ static constexpr bool value = type::value;
+};
} // namespace detail
// SYCL 2020 4.6.5. Queue class.
@@ -139,7 +163,82 @@ class _LIBSYCL_EXPORT queue {
/// exceptions.
void wait();
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type.
+ ///
+ /// \param kernelFunc is the kernel functor or lambda.
+ /// \return an event that represents the status of the submitted kernel.
+ template <typename KernelName, typename KernelType>
+ event single_task(const KernelType &kernelFunc) {
+ return single_task<KernelName, KernelType>({}, kernelFunc);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type.
+ ///
+ /// \param depEvent is an event that specifies the kernel dependency.
+ /// \param kernelFunc is the kernel functor or lambda.
+ /// \return an event that represents the status of the submitted kernel.
+ template <typename KernelName, typename KernelType>
+ event single_task(event depEvent, const KernelType &kernelFunc) {
+ return single_task<KernelName, KernelType>({depEvent}, kernelFunc);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type.
+ ///
+ /// \param depEvents is a collection of events that specify the kernel
+ /// dependencies.
+ /// \param kernelFunc is the kernel functor or lambda.
+ /// \return an event that represents the status of the submitted kernel.
+ template <typename KernelName, typename KernelType>
+ event single_task(const std::vector<event> &depEvents,
+ const KernelType &kernelFunc) {
+ static_assert(
+ detail::CheckFunctionSignature<std::remove_reference_t<KernelType>,
+ void()>::value,
+ "sycl::queue::single_task() requires a kernel instead of a command "
+ "group");
+
+ setKernelParameters(depEvents);
+ submitSingleTask<KernelName, KernelType>(kernelFunc);
+ return getLastEvent();
+ }
+
private:
+ // Name of this function is defined by compiler. It generates call to this
+ // function in the host implementation of KernelFunc in submitSingleTask.
+ template <typename KN, typename... Args>
+ void sycl_kernel_launch(const char *KernelName, Args &&...args) {
+ static_assert(
+ sizeof...(args) == 1,
+ "sycl_kernel_launch expects only 2 arguments now: name of kernel and "
+ "callable object passed to kernel invocation by the user.");
+
+ auto FirstArg = std::get<0>(std::tie(args...));
+ submitKernelImpl(detail::getDeviceKernelInfo<KN>(KernelName), &FirstArg,
+ sizeof(FirstArg));
+ }
+
+#ifdef SYCL_LANGUAGE_VERSION
+# define _LIBSYCL_ENTRY_POINT_ATTR__(KernelName) \
+ [[clang::sycl_kernel_entry_point(KernelName)]]
+#else
+# define _LIBSYCL_ENTRY_POINT_ATTR__(KernelName)
+#endif // SYCL_LANGUAGE_VERSION
+
+ template <typename KernelName, typename KernelType>
+ _LIBSYCL_ENTRY_POINT_ATTR__(KernelName)
+ void submitSingleTask(const KernelType &KernelFunc) {
+ KernelFunc();
+ }
+
+ event getLastEvent();
+ void submitKernelImpl(detail::DeviceKernelInfo &KernelInfo, void *ArgData,
+ size_t ArgSize);
+ void setKernelParameters(const std::vector<event> &Events,
+ const detail::UnifiedRangeView &Range = {});
+
queue(const std::shared_ptr<detail::QueueImpl> &Impl) : impl(Impl) {}
std::shared_ptr<detail::QueueImpl> impl;
diff --git a/libsycl/src/detail/global_objects.cpp b/libsycl/src/detail/global_objects.cpp
index 35e32985e7cbb..fd94d772337d6 100644
--- a/libsycl/src/detail/global_objects.cpp
+++ b/libsycl/src/detail/global_objects.cpp
@@ -8,6 +8,7 @@
#include <detail/global_objects.hpp>
#include <detail/platform_impl.hpp>
+#include <detail/program_manager.hpp>
#ifdef _WIN32
# include <windows.h>
@@ -31,6 +32,7 @@ struct StaticVarShutdownHandler {
StaticVarShutdownHandler &
operator=(const StaticVarShutdownHandler &) = delete;
~StaticVarShutdownHandler() {
+ ProgramAndKernelManager::getInstance().releaseResources();
// No error reporting in shutdown
std::ignore = olShutDown();
}
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index 90d7c48d3d1c7..870d55198181b 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -18,6 +18,23 @@
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
+DeviceKernelInfo &_LIBSYCL_EXPORT
+getDeviceKernelInfo(std::string_view KernelName) {
+ return ProgramAndKernelManager::getInstance().getDeviceKernelInfo(KernelName);
+}
+
+DeviceKernelInfo &
+ProgramAndKernelManager::getDeviceKernelInfo(std::string_view KernelName) {
+ auto It = MDeviceKernelInfoMap.find(KernelName);
+ assert(It != MDeviceKernelInfoMap.end());
+ return It->second;
+}
+
+void ProgramAndKernelManager::releaseResources() {
+ MDeviceKernelInfoMap.clear();
+ MDeviceImageManagers.clear();
+}
+
static inline bool checkFatBinVersion(const __sycl_tgt_bin_desc &FatbinDesc) {
return FatbinDesc.Version == SupportedOffloadBinaryVersion;
}
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index f5d3a1c8dc6dd..da56f6ec9706d 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -81,6 +81,12 @@ class ProgramAndKernelManager {
ol_symbol_handle_t getOrCreateKernel(DeviceKernelInfo &KernelInfo,
DeviceImpl &Device);
+ /// \return kernel info for the kernel with the specified name.
+ DeviceKernelInfo &getDeviceKernelInfo(std::string_view KernelName);
+
+ /// Release device image managers and corresponding resources.
+ void releaseResources();
+
private:
ProgramAndKernelManager() = default;
~ProgramAndKernelManager() = default;
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index 1d378f0ab5ef9..ff194f4d4152b 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -10,11 +10,41 @@
#include <detail/device_impl.hpp>
#include <detail/event_impl.hpp>
+#include <detail/program_manager.hpp>
+
+#include <algorithm>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
+static void setKernelLaunchArgs(const detail::UnifiedRangeView &Range,
+ ol_kernel_launch_size_args_t &ArgsToSet) {
+ assert(Range.MDims < 4 && "Invalid dimensions.");
+ uint32_t GlobalSize[3] = {1, 1, 1};
+ if (Range.MGlobalSize) {
+ for (auto I = 0; I < Range.MDims; I++) {
+ GlobalSize[I] = static_cast<uint32_t>(Range.MGlobalSize[I]);
+ }
+ }
+
+ uint32_t GroupSize[3] = {1, 1, 1};
+ if (Range.MLocalSize) {
+ for (auto I = 0; I < Range.MDims; I++) {
+ GroupSize[I] = static_cast<uint32_t>(Range.MLocalSize[I]);
+ }
+ }
+
+ ArgsToSet.Dimensions = Range.MDims;
+ ArgsToSet.NumGroups.x = GlobalSize[0] / GroupSize[0];
+ ArgsToSet.NumGroups.y = GlobalSize[1] / GroupSize[1];
+ ArgsToSet.NumGroups.z = GlobalSize[2] / GroupSize[2];
+ ArgsToSet.GroupSize.x = GroupSize[0];
+ ArgsToSet.GroupSize.y = GroupSize[1];
+ ArgsToSet.GroupSize.z = GroupSize[2];
+ ArgsToSet.DynSharedMemory = 0;
+}
+
QueueImpl::QueueImpl(DeviceImpl &deviceImpl, const async_handler &asyncHandler,
const property_list &propList, PrivateTag)
: MIsInorder(false), MAsyncHandler(asyncHandler), MPropList(propList),
@@ -33,5 +63,86 @@ backend QueueImpl::getBackend() const noexcept { return MDevice.getBackend(); }
void QueueImpl::wait() { callAndThrow(olSyncQueue, MOffloadQueue); }
+static bool checkEventsPlatformMatch(std::vector<EventImplPtr> &Events,
+ const PlatformImpl &QueuePlatform) {
+ // liboffload limitation to olWaitEvents. We can't do any extra handling for
+ // cross context/platform events without host task support now.
+ // "The input events can be from any queue on any device provided by the
+ // same platform as `Queue`."
+ return std::all_of(Events.cbegin(), Events.cend(),
+ [&QueuePlatform](const EventImplPtr &Event) {
+ return &Event->getPlatformImpl() == &QueuePlatform;
+ });
+}
+
+void QueueImpl::setKernelParameters(std::vector<EventImplPtr> &&Events,
+ const detail::UnifiedRangeView &Range) {
+ if (!checkEventsPlatformMatch(Events, MDevice.getPlatformImpl()))
+ throw sycl::exception(
+ sycl::make_error_code(sycl::errc::feature_not_supported),
+ "libsycl doesn't support cross-context/platform event dependencies "
+ "now.");
+
+ // TODO: this conversion and storing of only offload events is possible only
+ // while we don't have host tasks (or features based on host tasks, like
+ // streams). With them - it is very likely we should copy EventImplPtr
+ // (shared_ptr) and keep it here. Although it may differ if host tasks will be
+ // implemented on offload level (no data now).
+ assert(MCurrentSubmitInfo.DepEvents.empty() &&
+ "Kernel submission must clean up dependencies.");
+ MCurrentSubmitInfo.DepEvents.reserve(Events.size());
+ for (auto &Event : Events) {
+ assert(Event && "Event impl object can't be nullptr");
+ MCurrentSubmitInfo.DepEvents.push_back(Event->getHandle());
+ }
+ setKernelLaunchArgs(Range, MCurrentSubmitInfo.Range);
+}
+
+void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
+ size_t ArgSize) {
+ ol_symbol_handle_t Kernel =
+ detail::ProgramAndKernelManager::getInstance().getOrCreateKernel(
+ KernelInfo, MDevice);
+ assert(Kernel);
+
+ // TODO: liboffload supports only in-order queues and no cross context waiting
+ // is available now that means that this code is excessive but correct. I
+ // don't want to skip it and rely on default liboffload behaviour that is
+ // applicable for in-order queue only. Once OOO queues are added this waiting
+ // must be disabled for in-order queues. Once host tasks are added - cross
+ // context dependencies should be enabled and checked as well.
+ if (!MCurrentSubmitInfo.DepEvents.empty()) {
+ callAndThrow(olWaitEvents, MOffloadQueue,
+ MCurrentSubmitInfo.DepEvents.data(),
+ MCurrentSubmitInfo.DepEvents.size());
+ }
+
+ assert(ArgData && "At least one argument must exist");
+ assert(ArgSize && "Arguments size must be greater than 0");
+
+ // ol_kernel_launch_prop_t Props[2];
+ // Props[0].type = OL_KERNEL_LAUNCH_PROP_TYPE_SIZE;
+ // Props[0].data = &ArgSize;
+ // Props[1] = OL_KERNEL_LAUNCH_PROP_END;
+ auto Result =
+ olLaunchKernel(MOffloadQueue, MDevice.getOLHandle(), Kernel, &ArgData,
+ ArgSize, &MCurrentSubmitInfo.Range /*, Props*/);
+ // Clean up current kernel submit data to prepare structures for next
+ // submission.
+ MCurrentSubmitInfo.DepEvents.clear();
+ MCurrentSubmitInfo.Range = {};
+ if (isFailed(Result))
+ throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+ std::string("Kernel submission (") +
+ KernelInfo.getName().data() + ") failed with " +
+ formatCodeString(Result));
+
+ ol_event_handle_t NewEvent{};
+ callAndThrow(olCreateEvent, MOffloadQueue, &NewEvent);
+
+ MCurrentSubmitInfo.LastEvent =
+ EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl());
+}
+
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/queue_impl.hpp b/libsycl/src/detail/queue_impl.hpp
index 047cb121150f3..cda6ae2961c19 100644
--- a/libsycl/src/detail/queue_impl.hpp
+++ b/libsycl/src/detail/queue_impl.hpp
@@ -29,6 +29,8 @@ class ContextImpl;
class DeviceImpl;
class EventImpl;
+using EventImplPtr = std::shared_ptr<EventImpl>;
+
class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
struct PrivateTag {
explicit PrivateTag() = default;
@@ -69,13 +71,50 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
/// Waits for completion of all commands submitted to this queue.
void wait();
+ /// Enqueues a kernel to liboffload.
+ /// Kernel parameters like dependencies and range must be passed in advance by
+ /// calling setKernelParameters.
+ /// \param KernelInfo a kernel info that is uniform between different
+ /// submissions of the same kernel.
+ /// \param TypelessArgs data about kernel arguments to be used for enqueue.
+ void submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
+ size_t ArgSize);
+
+ /// \return an event impl object that corresponds to the last kernel
+ /// submission in the calling thread.
+ EventImplPtr getLastEvent() {
+ assert(MCurrentSubmitInfo.LastEvent &&
+ "getLastEvent must be called after enqueue");
+ return MCurrentSubmitInfo.LastEvent;
+ }
+
+ /// Sets kernel parameters to be used in the next submitKernelImpl call.
+ /// Must be called prior to a submitKernelImpl call.
+ /// \param Events a collection of events that the kernal depends on.
+ /// \param Range a unified range view of the execution range.
+ void setKernelParameters(std::vector<EventImplPtr> &&Events,
+ const detail::UnifiedRangeView &Range);
+
private:
+ // Queue features.
ol_queue_handle_t MOffloadQueue = {};
const bool MIsInorder;
const async_handler MAsyncHandler;
const property_list MPropList;
DeviceImpl &MDevice;
ContextImpl &MContext;
+
+ // Submit data.
+ struct KernelSubmitInfo {
+ EventImplPtr LastEvent;
+ ol_kernel_launch_size_args_t Range;
+ // TODO: consider storing EventImplPtr here, it will work with plain handle
+ // only because submission is done within queue::submit call. Otherwise we
+ // need to ensure that event handle is still alive by keeping our own copy
+ // of EventImpl.
+ std::vector<ol_event_handle_t> DepEvents;
+ };
+ inline static thread_local KernelSubmitInfo MCurrentSubmitInfo = {};
};
} // namespace detail
diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp
index 6584a6e080ec3..b57324219e46b 100644
--- a/libsycl/src/queue.cpp
+++ b/libsycl/src/queue.cpp
@@ -35,4 +35,23 @@ bool queue::is_in_order() const { return impl->isInOrder(); }
void queue::wait() { impl->wait(); }
+event queue::getLastEvent() {
+ return detail::createSyclObjFromImpl<event>(impl->getLastEvent());
+}
+
+void queue::setKernelParameters(const std::vector<event> &Events,
+ const detail::UnifiedRangeView &Range) {
+ std::vector<detail::EventImplPtr> DepEventImplRefs;
+ DepEventImplRefs.reserve(Events.size());
+ for (const auto &Event : Events) {
+ DepEventImplRefs.push_back(detail::getSyclObjImpl(Event));
+ }
+ return impl->setKernelParameters(std::move(DepEventImplRefs), Range);
+}
+
+void queue::submitKernelImpl(detail::DeviceKernelInfo &KernelInfo,
+ void *ArgData, size_t ArgSize) {
+ impl->submitKernelImpl(KernelInfo, ArgData, ArgSize);
+}
+
_LIBSYCL_END_NAMESPACE_SYCL
More information about the llvm-commits
mailing list