[llvm] [offload][sycl] add context parameter to olCreateProgram (PR #218387)
Łukasz Plewa via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 31 04:35:54 PDT 2026
https://github.com/lplewa updated https://github.com/llvm/llvm-project/pull/218387
>From 8289762f13657cb82b97cb15f7fa7cea2efea9ee Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?=C5=81ukasz=20Plewa?= <lukasz.plewa at intel.com>
Date: Wed, 19 Aug 2026 17:53:25 +0200
Subject: [PATCH 1/3] [offload][sycl] add context parameter to olCreateProgram
This patch is the 3rd patch in the context patch series.
This change is relatively simple compared to the others: we just introduce
context to the create program API and pass it down through the plugin
interface to the plugins.
---
libsycl/src/detail/device_image_wrapper.cpp | 12 +++--
libsycl/src/detail/device_image_wrapper.hpp | 9 +++-
libsycl/src/detail/program_manager.cpp | 3 +-
libsycl/src/detail/program_manager.hpp | 2 +
libsycl/src/detail/queue_impl.cpp | 2 +-
libsycl/unittests/mock/helpers.cpp | 5 +-
libsycl/unittests/mock/helpers.hpp | 5 +-
libsycl/unittests/mock/mock.cpp | 5 +-
.../tools/llvm-gpu-loader/llvm-gpu-loader.cpp | 2 +-
llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h | 3 +-
.../kernel/src/LanguageRegistration.cpp | 3 +-
offload/liboffload/API/Program.td | 8 ++-
offload/liboffload/src/OffloadImpl.cpp | 18 ++++---
offload/plugins-nextgen/amdgpu/src/rtl.cpp | 4 +-
.../common/include/PluginInterface.h | 9 ++--
.../common/src/PluginInterface.cpp | 10 ++--
offload/plugins-nextgen/cuda/src/rtl.cpp | 4 +-
offload/plugins-nextgen/host/src/rtl.cpp | 4 +-
.../level_zero/include/L0Device.h | 4 +-
.../level_zero/include/L0Program.h | 8 ++-
.../level_zero/src/L0Device.cpp | 6 ++-
.../level_zero/src/L0Program.cpp | 2 +-
.../Conformance/lib/DeviceContext.cpp | 2 +-
.../unittests/OffloadAPI/common/Fixtures.hpp | 5 +-
.../event/olGetEventElapsedTime.cpp | 2 +-
.../OffloadAPI/program/olCreateProgram.cpp | 50 ++++++++++++++-----
.../OffloadAPI/symbol/olGetSymbol.cpp | 2 +-
27 files changed, 127 insertions(+), 62 deletions(-)
diff --git a/libsycl/src/detail/device_image_wrapper.cpp b/libsycl/src/detail/device_image_wrapper.cpp
index d5cb0135c2854..17a830614bd0a 100644
--- a/libsycl/src/detail/device_image_wrapper.cpp
+++ b/libsycl/src/detail/device_image_wrapper.cpp
@@ -13,12 +13,15 @@
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
-ProgramWrapper::ProgramWrapper(ol_device_handle_t Device,
+ProgramWrapper::ProgramWrapper(ol_context_handle_t Context,
+ ol_device_handle_t Device,
DeviceImageManager &DevImage) {
+ assert(Context);
assert(Device);
llvm::StringRef Image = DevImage.getOffloadBinary().getImage();
- callAndThrow(olCreateProgram, Device, Image.data(), Image.size(), &MProgram);
+ callAndThrow(olCreateProgram, Context, Device, Image.data(), Image.size(),
+ &MProgram);
}
ProgramWrapper::~ProgramWrapper() {
@@ -28,10 +31,11 @@ ProgramWrapper::~ProgramWrapper() {
}
ol_program_handle_t
-DeviceImageManager::getOrCreateProgram(ol_device_handle_t DeviceHandle) {
+DeviceImageManager::getOrCreateProgram(ol_context_handle_t ContextHandle,
+ 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));
+ std::forward_as_tuple(ContextHandle, DeviceHandle, *this));
return Iterator->second.getOLHandle();
}
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
index bb0957002d251..5dfd7d05ee6b2 100644
--- a/libsycl/src/detail/device_image_wrapper.hpp
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -35,11 +35,13 @@ class ProgramWrapper {
/// Constructs ProgramWrapper by creating a liboffload program with the
/// provided arguments.
///
+ /// \param Context is the context to use for program creation.
/// \param Device is the device to use for program creation.
/// \param DevImage is the 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);
+ ProgramWrapper(ol_context_handle_t Context, ol_device_handle_t Device,
+ DeviceImageManager &DevImage);
/// Releases the corresponding liboffload program handle by calling
/// olDestroyProgram.
@@ -78,11 +80,14 @@ class DeviceImageManager {
/// 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 ContextHandle the liboffload handle of the context to create the
+ /// program in.
/// \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);
+ ol_program_handle_t getOrCreateProgram(ol_context_handle_t ContextHandle,
+ ol_device_handle_t DeviceHandle);
protected:
std::unordered_map<ol_device_handle_t, ProgramWrapper> MPrograms;
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index 688b49f847e10..d0ca9ba9dfce2 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -129,6 +129,7 @@ static bool isImageCompatible(const DeviceImageManager &Image,
ol_symbol_handle_t
ProgramAndKernelManager::getOrCreateKernel(DeviceKernelInfo &KernelInfo,
+ ol_context_handle_t Context,
DeviceImpl &Device) {
std::lock_guard<std::mutex> KernelGuard(MDataCollectionMutex);
@@ -144,7 +145,7 @@ ProgramAndKernelManager::getOrCreateKernel(DeviceKernelInfo &KernelInfo,
KernelInfo.getName().data() + " was found");
auto DeviceHandle = Device.getOLHandle();
- auto Program = DeviceImage.getOrCreateProgram(DeviceHandle);
+ auto Program = DeviceImage.getOrCreateProgram(Context, DeviceHandle);
ol_symbol_handle_t Kernel{};
callAndThrow(olGetSymbol, Program, KernelInfo.getName().data(),
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index cf978796a0cfd..ebea47ddd55df 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -80,10 +80,12 @@ class ProgramAndKernelManager {
/// This method is thread-safe.
/// \param KernelInfo a set of kernel specific data: name, corresponding
/// device image, etc.
+ /// \param Context the context in which the underlying program is created.
/// \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,
+ ol_context_handle_t Context,
DeviceImpl &Device);
/// \return kernel info for the kernel with the specified name.
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index e0700129c0fa5..78be58ec71e99 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -130,7 +130,7 @@ void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
size_t ArgSize) {
ol_symbol_handle_t Kernel =
detail::ProgramAndKernelManager::getInstance().getOrCreateKernel(
- KernelInfo, MDevice);
+ KernelInfo, MContext.getOLHandleRef(), MDevice);
assert(Kernel);
handleEventDependencies(MCurrentSubmitInfo.DepEvents);
diff --git a/libsycl/unittests/mock/helpers.cpp b/libsycl/unittests/mock/helpers.cpp
index 86e656a19a0ea..45fe4c7f3aa52 100644
--- a/libsycl/unittests/mock/helpers.cpp
+++ b/libsycl/unittests/mock/helpers.cpp
@@ -182,9 +182,10 @@ void mock::MockLiboffload::initDefault() {
});
ON_CALL(*this, olCreateProgram)
- .WillByDefault([](ol_device_handle_t Device, const void *ProgData,
- size_t ProgDataSize,
+ .WillByDefault([](ol_context_handle_t Context, ol_device_handle_t Device,
+ const void *ProgData, size_t ProgDataSize,
ol_program_handle_t *Program) -> ol_result_t {
+ std::ignore = Context;
EXPECT_NE(Device, nullptr);
EXPECT_NE(ProgData, nullptr);
EXPECT_GT(ProgDataSize, 0);
diff --git a/libsycl/unittests/mock/helpers.hpp b/libsycl/unittests/mock/helpers.hpp
index 50ff09b2b02c8..64b3c2fc55f23 100644
--- a/libsycl/unittests/mock/helpers.hpp
+++ b/libsycl/unittests/mock/helpers.hpp
@@ -100,8 +100,9 @@ class MockLiboffload {
MOCK_METHOD(ol_result_t, olSyncQueue, (ol_queue_handle_t Queue));
MOCK_METHOD(ol_result_t, olDestroyEvent, (ol_event_handle_t Event));
MOCK_METHOD(ol_result_t, olCreateProgram,
- (ol_device_handle_t Device, const void *ProgData,
- size_t ProgDataSize, ol_program_handle_t *Program));
+ (ol_context_handle_t Context, ol_device_handle_t Device,
+ const void *ProgData, size_t ProgDataSize,
+ ol_program_handle_t *Program));
MOCK_METHOD(ol_result_t, olGetSymbol,
(ol_program_handle_t Program, const char *Name,
diff --git a/libsycl/unittests/mock/mock.cpp b/libsycl/unittests/mock/mock.cpp
index 551c14e568701..74131c4692874 100644
--- a/libsycl/unittests/mock/mock.cpp
+++ b/libsycl/unittests/mock/mock.cpp
@@ -73,9 +73,10 @@ ol_result_t olSyncQueue(ol_queue_handle_t Queue) {
return mock::getMockLiboffload().olSyncQueue(Queue);
}
-ol_result_t olCreateProgram(ol_device_handle_t Device, const void *ProgData,
+ol_result_t olCreateProgram(ol_context_handle_t Context,
+ ol_device_handle_t Device, const void *ProgData,
size_t ProgDataSize, ol_program_handle_t *Program) {
- return mock::getMockLiboffload().olCreateProgram(Device, ProgData,
+ return mock::getMockLiboffload().olCreateProgram(Context, Device, ProgData,
ProgDataSize, Program);
}
diff --git a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
index 92dc25e4a6d38..55b59faeefdb4 100644
--- a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
+++ b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
@@ -256,7 +256,7 @@ int main(int argc, const char **argv, const char **envp) {
OFFLOAD_ERR(olCreateContext(1, &Device, &Context));
ol_program_handle_t Program;
- OFFLOAD_ERR(olCreateProgram(Device, Image.getBufferStart(),
+ OFFLOAD_ERR(olCreateProgram(Context, Device, Image.getBufferStart(),
Image.getBufferSize(), &Program));
ol_queue_handle_t Queue;
diff --git a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h
index c257f2ea391b2..c1dc25cc76cbc 100644
--- a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h
+++ b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h
@@ -123,7 +123,8 @@ ol_result_t (*olIterateDevices)(ol_device_iterate_cb_t Callback,
ol_result_t (*olIsValidBinary)(ol_device_handle_t Device, const void *ProgData,
size_t ProgDataSize, bool *Valid);
-ol_result_t (*olCreateProgram)(ol_device_handle_t Device, const void *ProgData,
+ol_result_t (*olCreateProgram)(ol_context_handle_t Context,
+ ol_device_handle_t Device, const void *ProgData,
size_t ProgDataSize,
ol_program_handle_t *Program);
diff --git a/offload/languages/kernel/src/LanguageRegistration.cpp b/offload/languages/kernel/src/LanguageRegistration.cpp
index 77aecdff21f23..47decd3f369d4 100644
--- a/offload/languages/kernel/src/LanguageRegistration.cpp
+++ b/offload/languages/kernel/src/LanguageRegistration.cpp
@@ -80,6 +80,7 @@ void __tgt_register_lib(__tgt_bin_desc *Desc) {
RuntimeState &State = RuntimeState::get();
ThreadState &Thread = ThreadState::get();
ol_device_handle_t Device = Thread.getDefaultDevice();
+ ol_context_handle_t Context = State.getContext();
for (int32_t I = 0, E = Desc->NumDeviceImages; I < E; ++I) {
ol_program_handle_t Program = nullptr;
@@ -89,7 +90,7 @@ void __tgt_register_lib(__tgt_bin_desc *Desc) {
size_t ProgramSize =
(char *)DeviceImage.ImageEnd - (char *)DeviceImage.ImageStart;
ol_result_t Result =
- olCreateProgram(Device, ProgramData, ProgramSize, &Program);
+ olCreateProgram(Context, Device, ProgramData, ProgramSize, &Program);
if (Result && Result->Code) {
fprintf(stderr, "Failed to register device code (%i): %s\n", Result->Code,
diff --git a/offload/liboffload/API/Program.td b/offload/liboffload/API/Program.td
index 89b9dffbe0a48..ecc8fb73dad3f 100644
--- a/offload/liboffload/API/Program.td
+++ b/offload/liboffload/API/Program.td
@@ -11,11 +11,14 @@
//===----------------------------------------------------------------------===//
def olCreateProgram : Function {
- let desc = "Create a program for the device from the binary image pointed to by `ProgData`.";
+ let desc = "Create a program for the device from the binary image pointed to by `ProgData` within the given context.";
let details = [
"The provided `ProgData` will be copied and need not outlive the returned handle",
+ "The program is scoped to `Context` and `Device` must belong to it.",
+ "The program can only be used with queues that were created in the same context."
];
let params = [
+ Param<"ol_context_handle_t", "Context", "handle of the context", PARAM_IN>,
Param<"ol_device_handle_t", "Device", "handle of the device", PARAM_IN>,
Param<"const void*", "ProgData", "pointer to the program binary data", PARAM_IN>,
Param<"size_t", "ProgDataSize", "size of the program binary in bytes", PARAM_IN>,
@@ -25,6 +28,9 @@ def olCreateProgram : Function {
Return<"OL_ERRC_INVALID_BINARY", [
"If the buffer described by `ProgData` and `ProgDataSize` is not a valid binary image for the platform."
]>,
+ Return<"OL_ERRC_INVALID_DEVICE", [
+ "Device does not belong to `Context`"
+ ]>,
];
}
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index e59fed4b30c34..e666be7a5b287 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -133,9 +133,10 @@ struct ol_event_impl_t {
};
struct ol_program_impl_t {
- ol_program_impl_t(plugin::DeviceImageTy *Image,
+ ol_program_impl_t(ol_context_handle_t Context, plugin::DeviceImageTy *Image,
llvm::MemoryBufferRef DeviceImage)
- : Image(Image), DeviceImage(DeviceImage) {}
+ : Context(Context), Image(Image), DeviceImage(DeviceImage) {}
+ ol_context_handle_t Context;
plugin::DeviceImageTy *Image;
std::mutex SymbolListMutex;
llvm::MemoryBufferRef DeviceImage;
@@ -1166,16 +1167,21 @@ Error olMemPrefetch_impl(ol_queue_handle_t Queue, size_t Count,
Queue->AsyncInfo);
}
-Error olCreateProgram_impl(ol_device_handle_t Device, const void *ProgData,
+Error olCreateProgram_impl(ol_context_handle_t Context,
+ ol_device_handle_t Device, const void *ProgData,
size_t ProgDataSize, ol_program_handle_t *Program) {
+ if (!Context->contains(Device))
+ return createOffloadError(ErrorCode::INVALID_DEVICE,
+ "device does not belong to the given context");
+
StringRef Buffer(reinterpret_cast<const char *>(ProgData), ProgDataSize);
- Expected<plugin::DeviceImageTy *> Res =
- Device->Device->loadBinary(Device->Device->Plugin, Buffer);
+ Expected<plugin::DeviceImageTy *> Res = Device->Device->loadBinary(
+ Device->Device->Plugin, Buffer, Context->PluginCtx.get());
if (!Res)
return Res.takeError();
assert(*Res && "loadBinary returned nullptr");
- *Program = new ol_program_impl_t(*Res, (*Res)->getMemoryBuffer());
+ *Program = new ol_program_impl_t(Context, *Res, (*Res)->getMemoryBuffer());
return Error::success();
}
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index 2ba4d80978f4a..281b9e3795a54 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -2686,8 +2686,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
/// Load the binary image into the device and allocate an image object.
Expected<DeviceImageTy *>
- loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
- int32_t ImageId) override {
+ loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, int32_t ImageId,
+ PluginContextTy * /*Context*/) override {
// Allocate and initialize the image object.
AMDGPUDeviceImageTy *AMDImage = Plugin.allocate<AMDGPUDeviceImageTy>();
new (AMDImage) AMDGPUDeviceImageTy(ImageId, *this, std::move(TgtImage));
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 2f8d52c830a64..7b7a62bc2d467 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -983,11 +983,14 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
Error deinit(GenericPluginTy &Plugin);
virtual Error deinitImpl() = 0;
- /// Load the binary image into the device and return the target table.
+ /// Load the binary image into the device and return the target table. When
+ /// \p Context is null the plugin's driver-scoped default context is used.
Expected<DeviceImageTy *> loadBinary(GenericPluginTy &Plugin,
- StringRef TgtImage);
+ StringRef TgtImage,
+ PluginContextTy *Context);
virtual Expected<DeviceImageTy *>
- loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, int32_t ImageId) = 0;
+ loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, int32_t ImageId,
+ PluginContextTy *Context) = 0;
/// Unload a previously loaded Image from the device
Error unloadBinary(DeviceImageTy *Image);
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 9d6135da6d470..573369dc3ba57 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -645,8 +645,9 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
return deinitImpl();
}
-Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,
- StringRef InputTgtImage) {
+Expected<DeviceImageTy *>
+GenericDeviceTy::loadBinary(GenericPluginTy &Plugin, StringRef InputTgtImage,
+ PluginContextTy *Context) {
ODBG(OLDT_Init) << "Load data from image "
<< static_cast<const void *>(InputTgtImage.bytes_begin());
@@ -674,7 +675,8 @@ Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,
// Load the binary and allocate the image object. Use the next available id
// for the image id, which is the number of previously loaded images.
- auto ImageOrErr = loadBinaryImpl(std::move(Buffer), LoadedImages.size());
+ auto ImageOrErr =
+ loadBinaryImpl(std::move(Buffer), LoadedImages.size(), Context);
if (!ImageOrErr)
return ImageOrErr.takeError();
DeviceImageTy *Image = *ImageOrErr;
@@ -1542,7 +1544,7 @@ int32_t GenericPluginTy::load_binary(int32_t DeviceId,
StringRef Buffer(reinterpret_cast<const char *>(TgtImage->ImageStart),
utils::getPtrDiff(TgtImage->ImageEnd, TgtImage->ImageStart));
- auto ImageOrErr = Device.loadBinary(*this, Buffer);
+ auto ImageOrErr = Device.loadBinary(*this, Buffer, /*Context=*/nullptr);
if (!ImageOrErr) {
auto Err = ImageOrErr.takeError();
REPORT() << "Failure to load binary image " << TgtImage << " on device "
diff --git a/offload/plugins-nextgen/cuda/src/rtl.cpp b/offload/plugins-nextgen/cuda/src/rtl.cpp
index 0f666ffb65e8f..72e5dcf115fe9 100644
--- a/offload/plugins-nextgen/cuda/src/rtl.cpp
+++ b/offload/plugins-nextgen/cuda/src/rtl.cpp
@@ -561,8 +561,8 @@ struct CUDADeviceTy : public GenericDeviceTy {
/// Load the binary image into the device and allocate an image object.
Expected<DeviceImageTy *>
- loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
- int32_t ImageId) override {
+ loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, int32_t ImageId,
+ PluginContextTy * /*Context*/) override {
if (auto Err = setContext())
return std::move(Err);
diff --git a/offload/plugins-nextgen/host/src/rtl.cpp b/offload/plugins-nextgen/host/src/rtl.cpp
index 3c5c29545d89d..55ada2f82c360 100644
--- a/offload/plugins-nextgen/host/src/rtl.cpp
+++ b/offload/plugins-nextgen/host/src/rtl.cpp
@@ -182,8 +182,8 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
/// Load the binary image into the device and allocate an image object.
Expected<DeviceImageTy *>
- loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
- int32_t ImageId) override {
+ loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, int32_t ImageId,
+ PluginContextTy * /*Context*/) override {
// Allocate and initialize the image object.
GenELF64DeviceImageTy *Image = Plugin.allocate<GenELF64DeviceImageTy>();
new (Image) GenELF64DeviceImageTy(ImageId, *this, std::move(TgtImage));
diff --git a/offload/plugins-nextgen/level_zero/include/L0Device.h b/offload/plugins-nextgen/level_zero/include/L0Device.h
index 84df2a2140446..d535b8abb0fc0 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Device.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Device.h
@@ -510,8 +510,8 @@ class L0DeviceTy final : public GenericDeviceTy {
// Generic device interface implementation.
Expected<DeviceImageTy *>
- loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
- int32_t ImageId) override;
+ loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, int32_t ImageId,
+ PluginContextTy *UserCtx) override;
Error unloadBinaryImpl(DeviceImageTy *Image) override;
Expected<void *> allocate(size_t Size, void *HstPtr, TargetAllocTy Kind,
size_t Alignment) override;
diff --git a/offload/plugins-nextgen/level_zero/include/L0Program.h b/offload/plugins-nextgen/level_zero/include/L0Program.h
index 68e901e984d22..e9fb2e778e0c6 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Program.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Program.h
@@ -21,6 +21,8 @@ class L0DeviceTy;
class L0ProgramBuilderTy {
L0DeviceTy &Device;
+ /// L0 context that owns the built modules.
+ ze_context_handle_t ZeContext;
std::unique_ptr<MemoryBuffer> Image;
/// Handle multiple modules within a single target image.
llvm::SmallVector<ze_module_handle_t> Modules;
@@ -39,11 +41,13 @@ class L0ProgramBuilderTy {
Error linkModules();
public:
- L0ProgramBuilderTy(L0DeviceTy &Device, std::unique_ptr<MemoryBuffer> &&Image)
- : Device(Device), Image(std::move(Image)) {}
+ L0ProgramBuilderTy(L0DeviceTy &Device, ze_context_handle_t ZeContext,
+ std::unique_ptr<MemoryBuffer> &&Image)
+ : Device(Device), ZeContext(ZeContext), Image(std::move(Image)) {}
~L0ProgramBuilderTy() = default;
L0DeviceTy &getL0Device() const { return Device; }
+ ze_context_handle_t getZeContext() const { return ZeContext; }
ze_module_handle_t getGlobalModule() const { return GlobalModule; }
llvm::SmallVector<ze_module_handle_t> &getModules() { return Modules; }
diff --git a/offload/plugins-nextgen/level_zero/src/L0Device.cpp b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
index 08c45b534e16c..b07aa91ff44e8 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Device.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
@@ -213,7 +213,7 @@ Error L0DeviceTy::deinitImpl() {
Expected<DeviceImageTy *>
L0DeviceTy::loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
- int32_t ImageId) {
+ int32_t ImageId, PluginContextTy *UserCtx) {
auto *PGM = getProgramFromImage(TgtImage->getMemBufferRef());
if (PGM) {
// Program already exists.
@@ -234,7 +234,9 @@ L0DeviceTy::loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
CompilationOptions += " ";
CompilationOptions += Options.InternalCompilationOptions;
- L0ProgramBuilderTy Builder(*this, std::move(TgtImage));
+ auto &Ctx = UserCtx ? static_cast<LevelZeroPluginContextTy &>(*UserCtx)
+ : L0Context.getDefaultUserCtx();
+ L0ProgramBuilderTy Builder(*this, Ctx.getZeContext(), std::move(TgtImage));
if (auto Err = Builder.buildModules(CompilationOptions))
return std::move(Err);
diff --git a/offload/plugins-nextgen/level_zero/src/L0Program.cpp b/offload/plugins-nextgen/level_zero/src/L0Program.cpp
index 6cadfe44d0813..7786164ae0040 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Program.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Program.cpp
@@ -94,7 +94,7 @@ Error L0ProgramBuilderTy::addModule(size_t Size, const uint8_t *Image,
ModuleDesc.pBuildFlags = BuildOptions.c_str();
ModuleDesc.pConstants = &SpecConstants;
ze_result_t RC;
- CALL_ZE(RC, zeModuleCreate, L0Device.getZeContext(), L0Device.getZeDevice(),
+ CALL_ZE(RC, zeModuleCreate, getZeContext(), L0Device.getZeDevice(),
&ModuleDesc, &Module, &BuildLog);
if (BuildLog)
zeModuleBuildLogDestroy(BuildLog);
diff --git a/offload/unittests/Conformance/lib/DeviceContext.cpp b/offload/unittests/Conformance/lib/DeviceContext.cpp
index f3f0f8f394323..d81f5f0c5867a 100644
--- a/offload/unittests/Conformance/lib/DeviceContext.cpp
+++ b/offload/unittests/Conformance/lib/DeviceContext.cpp
@@ -252,7 +252,7 @@ DeviceContext::loadBinary(llvm::StringRef Directory,
ol_program_handle_t ProgramHandle = nullptr;
const ol_result_t OlResult =
- olCreateProgram(DeviceHandle, BinaryData->getBufferStart(),
+ olCreateProgram(Context, DeviceHandle, BinaryData->getBufferStart(),
BinaryData->getBufferSize(), &ProgramHandle);
if (OlResult != OL_SUCCESS) {
diff --git a/offload/unittests/OffloadAPI/common/Fixtures.hpp b/offload/unittests/OffloadAPI/common/Fixtures.hpp
index 9ccc3c75314ee..a05be01648ebc 100644
--- a/offload/unittests/OffloadAPI/common/Fixtures.hpp
+++ b/offload/unittests/OffloadAPI/common/Fixtures.hpp
@@ -261,7 +261,8 @@ struct OffloadProgramTestWithParam : OffloadDeviceTestWithParam<T> {
ASSERT_TRUE(TestEnvironment::loadDeviceBinary(ProgramName, this->Device,
DeviceBin));
ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
- ASSERT_SUCCESS(olCreateProgram(this->Device, DeviceBin->getBufferStart(),
+ ASSERT_SUCCESS(olCreateProgram(this->Context, this->Device,
+ DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), &Program));
}
@@ -346,7 +347,7 @@ struct LaunchKernelTestBase : OffloadQueueTest {
RETURN_ON_FATAL_FAILURE(OffloadQueueTest::SetUp());
ASSERT_TRUE(TestEnvironment::loadDeviceBinary(program, Device, DeviceBin));
ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
- ASSERT_SUCCESS(olCreateProgram(Device, DeviceBin->getBufferStart(),
+ ASSERT_SUCCESS(olCreateProgram(Context, Device, DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), &Program));
LaunchArgs.Dimensions = 1;
diff --git a/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp b/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp
index 33a62921f4ae4..356f04bfd9bdb 100644
--- a/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp
+++ b/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp
@@ -19,7 +19,7 @@ struct olGetEventElapsedTimeTest : OffloadQueueTest {
SKIP_KNOWN_FAILURE(LevelZero{"unsupported feature"});
ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Device, DeviceBin));
- ASSERT_SUCCESS(olCreateProgram(Device, DeviceBin->getBufferStart(),
+ ASSERT_SUCCESS(olCreateProgram(Context, Device, DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), &Program));
ASSERT_SUCCESS(olGetSymbol(Program, "foo", OL_SYMBOL_KIND_KERNEL, &Kernel));
diff --git a/offload/unittests/OffloadAPI/program/olCreateProgram.cpp b/offload/unittests/OffloadAPI/program/olCreateProgram.cpp
index 59629c79caff4..13811af32d479 100644
--- a/offload/unittests/OffloadAPI/program/olCreateProgram.cpp
+++ b/offload/unittests/OffloadAPI/program/olCreateProgram.cpp
@@ -20,7 +20,7 @@ TEST_P(olCreateProgramTest, Success) {
ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
ol_program_handle_t Program;
- ASSERT_SUCCESS(olCreateProgram(Device, DeviceBin->getBufferStart(),
+ ASSERT_SUCCESS(olCreateProgram(Context, Device, DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), &Program));
ASSERT_NE(Program, nullptr);
@@ -34,7 +34,7 @@ TEST_P(olCreateProgramTest, JITSuccess) {
ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
ol_program_handle_t Program;
- ASSERT_SUCCESS(olCreateProgram(Device, DeviceBin->getBufferStart(),
+ ASSERT_SUCCESS(olCreateProgram(Context, Device, DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), &Program));
ASSERT_NE(Program, nullptr);
@@ -46,6 +46,18 @@ TEST_P(olCreateProgramTest, JITSuccess) {
ASSERT_SUCCESS(olDestroyProgram(Program));
}
+TEST_P(olCreateProgramTest, NullContextHandle) {
+
+ std::unique_ptr<llvm::MemoryBuffer> DeviceBin;
+ ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Device, DeviceBin));
+ ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
+
+ ol_program_handle_t Program;
+ ASSERT_ERROR(OL_ERRC_INVALID_NULL_HANDLE,
+ olCreateProgram(nullptr, Device, DeviceBin->getBufferStart(),
+ DeviceBin->getBufferSize(), &Program));
+}
+
TEST_P(olCreateProgramTest, NullDeviceHandle) {
std::unique_ptr<llvm::MemoryBuffer> DeviceBin;
@@ -54,7 +66,7 @@ TEST_P(olCreateProgramTest, NullDeviceHandle) {
ol_program_handle_t Program;
ASSERT_ERROR(OL_ERRC_INVALID_NULL_HANDLE,
- olCreateProgram(nullptr, DeviceBin->getBufferStart(),
+ olCreateProgram(Context, nullptr, DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), &Program));
}
@@ -65,9 +77,9 @@ TEST_P(olCreateProgramTest, NullProgData) {
ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
ol_program_handle_t Program;
- ASSERT_ERROR(
- OL_ERRC_INVALID_NULL_POINTER,
- olCreateProgram(Device, nullptr, DeviceBin->getBufferSize(), &Program));
+ ASSERT_ERROR(OL_ERRC_INVALID_NULL_POINTER,
+ olCreateProgram(Context, Device, nullptr,
+ DeviceBin->getBufferSize(), &Program));
}
TEST_P(olCreateProgramTest, NullOutputProgram) {
@@ -77,7 +89,7 @@ TEST_P(olCreateProgramTest, NullOutputProgram) {
ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
ASSERT_ERROR(OL_ERRC_INVALID_NULL_POINTER,
- olCreateProgram(Device, DeviceBin->getBufferStart(),
+ olCreateProgram(Context, Device, DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), nullptr));
}
@@ -88,9 +100,9 @@ TEST_P(olCreateProgramTest, ZeroSizeBinary) {
ol_program_handle_t Program = nullptr;
- ASSERT_ERROR(
- OL_ERRC_INVALID_BINARY,
- olCreateProgram(Device, DeviceBin->getBufferStart(), 0, &Program));
+ ASSERT_ERROR(OL_ERRC_INVALID_BINARY,
+ olCreateProgram(Context, Device, DeviceBin->getBufferStart(), 0,
+ &Program));
ASSERT_EQ(Program, nullptr);
}
@@ -99,8 +111,8 @@ TEST_P(olCreateProgramTest, InvalidBinary) {
ol_program_handle_t Program = nullptr;
ASSERT_ERROR(OL_ERRC_INVALID_BINARY,
- olCreateProgram(Device, InvalidBinary, sizeof(InvalidBinary) - 1,
- &Program));
+ olCreateProgram(Context, Device, InvalidBinary,
+ sizeof(InvalidBinary) - 1, &Program));
ASSERT_EQ(Program, nullptr);
}
@@ -120,7 +132,19 @@ TEST_P(olCreateProgramTest, WrongArchitecture) {
ol_program_handle_t Program = nullptr;
ASSERT_ERROR(OL_ERRC_INVALID_BINARY,
- olCreateProgram(Device, ForeignBin->getBufferStart(),
+ olCreateProgram(Context, Device, ForeignBin->getBufferStart(),
ForeignBin->getBufferSize(), &Program));
ASSERT_EQ(Program, nullptr);
}
+
+TEST_P(olCreateProgramTest, InvalidDeviceNotInContext) {
+ std::unique_ptr<llvm::MemoryBuffer> DeviceBin;
+ ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Device, DeviceBin));
+ ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
+
+ ol_program_handle_t Program = nullptr;
+ ASSERT_ERROR(OL_ERRC_INVALID_DEVICE,
+ olCreateProgram(Context, Host, DeviceBin->getBufferStart(),
+ DeviceBin->getBufferSize(), &Program));
+ ASSERT_EQ(Program, nullptr);
+}
diff --git a/offload/unittests/OffloadAPI/symbol/olGetSymbol.cpp b/offload/unittests/OffloadAPI/symbol/olGetSymbol.cpp
index 1f496b9c6e1ae..021aa35b89bc6 100644
--- a/offload/unittests/OffloadAPI/symbol/olGetSymbol.cpp
+++ b/offload/unittests/OffloadAPI/symbol/olGetSymbol.cpp
@@ -18,7 +18,7 @@ struct olGetSymbolGlobalTest : OffloadQueueTest {
RETURN_ON_FATAL_FAILURE(OffloadQueueTest::SetUp());
ASSERT_TRUE(TestEnvironment::loadDeviceBinary("global", Device, DeviceBin));
ASSERT_GE(DeviceBin->getBufferSize(), 0lu);
- ASSERT_SUCCESS(olCreateProgram(Device, DeviceBin->getBufferStart(),
+ ASSERT_SUCCESS(olCreateProgram(Context, Device, DeviceBin->getBufferStart(),
DeviceBin->getBufferSize(), &Program));
}
>From 920c14d1581463ff8e7468ef0e83ae4d3425c495 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?=C5=81ukasz=20Plewa?= <lukasz.plewa at intel.com>
Date: Wed, 26 Aug 2026 17:49:20 +0200
Subject: [PATCH 2/3] review fixes
---
libsycl/src/detail/program_manager.cpp | 10 +++++-----
libsycl/src/detail/program_manager.hpp | 6 ++++--
libsycl/src/detail/queue_impl.cpp | 2 +-
libsycl/unittests/mock/helpers.cpp | 2 +-
4 files changed, 11 insertions(+), 9 deletions(-)
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index d0ca9ba9dfce2..34e80435290e7 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -11,6 +11,7 @@
#include <sycl/__impl/detail/get_device_kernel_info.hpp>
#include <sycl/__impl/exception.hpp>
+#include <detail/context_impl.hpp>
#include <detail/device_impl.hpp>
#include <detail/offload/offload_utils.hpp>
@@ -127,10 +128,8 @@ static bool isImageCompatible(const DeviceImageManager &Image,
return IsValid;
}
-ol_symbol_handle_t
-ProgramAndKernelManager::getOrCreateKernel(DeviceKernelInfo &KernelInfo,
- ol_context_handle_t Context,
- DeviceImpl &Device) {
+ol_symbol_handle_t ProgramAndKernelManager::getOrCreateKernel(
+ DeviceKernelInfo &KernelInfo, ContextImpl &Context, DeviceImpl &Device) {
std::lock_guard<std::mutex> KernelGuard(MDataCollectionMutex);
@@ -145,7 +144,8 @@ ProgramAndKernelManager::getOrCreateKernel(DeviceKernelInfo &KernelInfo,
KernelInfo.getName().data() + " was found");
auto DeviceHandle = Device.getOLHandle();
- auto Program = DeviceImage.getOrCreateProgram(Context, DeviceHandle);
+ auto Program =
+ DeviceImage.getOrCreateProgram(Context.getOLHandleRef(), DeviceHandle);
ol_symbol_handle_t Kernel{};
callAndThrow(olGetSymbol, Program, KernelInfo.getName().data(),
diff --git a/libsycl/src/detail/program_manager.hpp b/libsycl/src/detail/program_manager.hpp
index ebea47ddd55df..59f77633debb1 100644
--- a/libsycl/src/detail/program_manager.hpp
+++ b/libsycl/src/detail/program_manager.hpp
@@ -54,6 +54,7 @@ _LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
+class ContextImpl;
class DeviceImpl;
/// A class to manage programs and kernels.
@@ -80,12 +81,13 @@ class ProgramAndKernelManager {
/// This method is thread-safe.
/// \param KernelInfo a set of kernel specific data: name, corresponding
/// device image, etc.
- /// \param Context the context in which the underlying program is created.
+ /// \param Context the context in which the underlying program must be
+ /// created.
/// \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,
- ol_context_handle_t Context,
+ ContextImpl &Context,
DeviceImpl &Device);
/// \return kernel info for the kernel with the specified name.
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index 78be58ec71e99..881e4617b3a85 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -130,7 +130,7 @@ void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
size_t ArgSize) {
ol_symbol_handle_t Kernel =
detail::ProgramAndKernelManager::getInstance().getOrCreateKernel(
- KernelInfo, MContext.getOLHandleRef(), MDevice);
+ KernelInfo, MContext, MDevice);
assert(Kernel);
handleEventDependencies(MCurrentSubmitInfo.DepEvents);
diff --git a/libsycl/unittests/mock/helpers.cpp b/libsycl/unittests/mock/helpers.cpp
index 45fe4c7f3aa52..b3838daa00742 100644
--- a/libsycl/unittests/mock/helpers.cpp
+++ b/libsycl/unittests/mock/helpers.cpp
@@ -185,7 +185,7 @@ void mock::MockLiboffload::initDefault() {
.WillByDefault([](ol_context_handle_t Context, ol_device_handle_t Device,
const void *ProgData, size_t ProgDataSize,
ol_program_handle_t *Program) -> ol_result_t {
- std::ignore = Context;
+ EXPECT_NE(Context, nullptr);
EXPECT_NE(Device, nullptr);
EXPECT_NE(ProgData, nullptr);
EXPECT_GT(ProgDataSize, 0);
>From d8e5f5f39ee5e6d8da34c13e39bcc6c5592ef922 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?=C5=81ukasz=20Plewa?= <lukasz.plewa at intel.com>
Date: Mon, 31 Aug 2026 12:11:43 +0200
Subject: [PATCH 3/3] review fixes - program cache
---
.../plugins-nextgen/level_zero/include/L0Device.h | 11 ++++++-----
.../plugins-nextgen/level_zero/include/L0Program.h | 12 ++++++++++--
offload/plugins-nextgen/level_zero/src/L0Device.cpp | 10 ++++++----
3 files changed, 22 insertions(+), 11 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/include/L0Device.h b/offload/plugins-nextgen/level_zero/include/L0Device.h
index d535b8abb0fc0..879c43acba596 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Device.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Device.h
@@ -213,9 +213,10 @@ class L0DeviceTy final : public GenericDeviceTy {
}
ze_module_handle_t *getGlobalModulesArray() { return GlobalModules.data(); }
- L0ProgramTy *getProgramFromImage(MemoryBufferRef Image) {
+ L0ProgramTy *getProgramFromImage(MemoryBufferRef Image,
+ ze_context_handle_t ZeContext) {
for (auto &PGM : Programs)
- if (PGM.getMemoryBuffer() == Image)
+ if (PGM.getMemoryBuffer() == Image && PGM.getZeContext() == ZeContext)
return &PGM;
return nullptr;
}
@@ -234,9 +235,9 @@ class L0DeviceTy final : public GenericDeviceTy {
auto ImageOrErr = Builder.getELF();
if (!ImageOrErr)
return ImageOrErr.takeError();
- Programs.emplace_back(ImageId, *this, std::move(*ImageOrErr),
- Builder.getGlobalModule(),
- std::move(Builder.getModules()));
+ Programs.emplace_back(
+ ImageId, *this, std::move(*ImageOrErr), Builder.getGlobalModule(),
+ std::move(Builder.getModules()), Builder.getZeContext());
return Programs.back();
}
diff --git a/offload/plugins-nextgen/level_zero/include/L0Program.h b/offload/plugins-nextgen/level_zero/include/L0Program.h
index e9fb2e778e0c6..5e29e4463e4db 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Program.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Program.h
@@ -74,6 +74,10 @@ class L0ProgramTy : public DeviceImageTy {
/// Module that contains global data including device RTL.
ze_module_handle_t GlobalModule = nullptr;
+ /// L0 context the modules were built against. Cached programs are only
+ /// reusable for lookups that share this context.
+ ze_context_handle_t ZeContext = nullptr;
+
L0DeviceTy &getL0Device() const;
public:
@@ -82,9 +86,11 @@ class L0ProgramTy : public DeviceImageTy {
L0ProgramTy(int32_t ImageId, GenericDeviceTy &Device,
std::unique_ptr<MemoryBuffer> Image,
ze_module_handle_t GlobalModule,
- llvm::SmallVector<ze_module_handle_t> &&Modules)
+ llvm::SmallVector<ze_module_handle_t> &&Modules,
+ ze_context_handle_t ZeContext)
: DeviceImageTy(ImageId, Device, std::move(Image)),
- Modules(std::move(Modules)), GlobalModule(GlobalModule) {}
+ Modules(std::move(Modules)), GlobalModule(GlobalModule),
+ ZeContext(ZeContext) {}
~L0ProgramTy() = default;
L0ProgramTy(const L0ProgramTy &Other) = delete;
@@ -94,6 +100,8 @@ class L0ProgramTy : public DeviceImageTy {
Error deinit();
+ ze_context_handle_t getZeContext() const { return ZeContext; }
+
static L0ProgramTy &makeL0Program(DeviceImageTy &Device) {
return static_cast<L0ProgramTy &>(Device);
}
diff --git a/offload/plugins-nextgen/level_zero/src/L0Device.cpp b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
index b07aa91ff44e8..0919dfc4a279b 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Device.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
@@ -214,7 +214,11 @@ Error L0DeviceTy::deinitImpl() {
Expected<DeviceImageTy *>
L0DeviceTy::loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
int32_t ImageId, PluginContextTy *UserCtx) {
- auto *PGM = getProgramFromImage(TgtImage->getMemBufferRef());
+ auto &Ctx = UserCtx ? static_cast<LevelZeroPluginContextTy &>(*UserCtx)
+ : L0Context.getDefaultUserCtx();
+ auto ZeContext = Ctx.getZeContext();
+
+ auto *PGM = getProgramFromImage(TgtImage->getMemBufferRef(), ZeContext);
if (PGM) {
// Program already exists.
return PGM;
@@ -234,9 +238,7 @@ L0DeviceTy::loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,
CompilationOptions += " ";
CompilationOptions += Options.InternalCompilationOptions;
- auto &Ctx = UserCtx ? static_cast<LevelZeroPluginContextTy &>(*UserCtx)
- : L0Context.getDefaultUserCtx();
- L0ProgramBuilderTy Builder(*this, Ctx.getZeContext(), std::move(TgtImage));
+ L0ProgramBuilderTy Builder(*this, ZeContext, std::move(TgtImage));
if (auto Err = Builder.buildModules(CompilationOptions))
return std::move(Err);
More information about the llvm-commits
mailing list