[llvm] [offload][sycl] add context parameter to olCreateProgram (PR #218387)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 24 05:27:12 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Łukasz Plewa (lplewa)
<details>
<summary>Changes</summary>
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.
---
Patch is 32.11 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/218387.diff
27 Files Affected:
- (modified) libsycl/src/detail/device_image_wrapper.cpp (+8-4)
- (modified) libsycl/src/detail/device_image_wrapper.hpp (+7-2)
- (modified) libsycl/src/detail/program_manager.cpp (+2-1)
- (modified) libsycl/src/detail/program_manager.hpp (+2)
- (modified) libsycl/src/detail/queue_impl.cpp (+1-1)
- (modified) libsycl/unittests/mock/helpers.cpp (+3-2)
- (modified) libsycl/unittests/mock/helpers.hpp (+3-2)
- (modified) libsycl/unittests/mock/mock.cpp (+3-2)
- (modified) llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp (+1-1)
- (modified) llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h (+2-1)
- (modified) offload/languages/kernel/src/LanguageRegistration.cpp (+2-1)
- (modified) offload/liboffload/API/Program.td (+7-1)
- (modified) offload/liboffload/src/OffloadImpl.cpp (+12-6)
- (modified) offload/plugins-nextgen/amdgpu/src/rtl.cpp (+2-2)
- (modified) offload/plugins-nextgen/common/include/PluginInterface.h (+6-3)
- (modified) offload/plugins-nextgen/common/src/PluginInterface.cpp (+5-3)
- (modified) offload/plugins-nextgen/cuda/src/rtl.cpp (+2-2)
- (modified) offload/plugins-nextgen/host/src/rtl.cpp (+2-2)
- (modified) offload/plugins-nextgen/level_zero/include/L0Device.h (+2-2)
- (modified) offload/plugins-nextgen/level_zero/include/L0Program.h (+6-2)
- (modified) offload/plugins-nextgen/level_zero/src/L0Device.cpp (+4-2)
- (modified) offload/plugins-nextgen/level_zero/src/L0Program.cpp (+1-1)
- (modified) offload/unittests/Conformance/lib/DeviceContext.cpp (+1-1)
- (modified) offload/unittests/OffloadAPI/common/Fixtures.hpp (+4-3)
- (modified) offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp (+1-1)
- (modified) offload/unittests/OffloadAPI/program/olCreateProgram.cpp (+37-13)
- (modified) offload/unittests/OffloadAPI/symbol/olGetSymbol.cpp (+1-1)
``````````diff
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 2898b39daafad..2ffc90d3ecdfb 100644
--- a/libsycl/unittests/mock/helpers.cpp
+++ b/libsycl/unittests/mock/helpers.cpp
@@ -186,9 +186,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 2219e16ea0365..bf88c7b200fb7 100644
--- a/offload/languages/kernel/src/LanguageRegistration.cpp
+++ b/offload/languages/kernel/src/LanguageRegistration.cpp
@@ -77,6 +77,7 @@ struct __tgt_bin_desc {
void __tgt_register_lib(__tgt_bin_desc *Desc) {
// TODO: For each device, lazily.
ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ ol_context_handle_t Context = RuntimeState::getContext();
for (int32_t I = 0, E = Desc->NumDeviceImages; I < E; ++I) {
ol_program_handle_t Program = nullptr;
@@ -86,7 +87,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 883d368ea2aab..c8a0ec9b958b9 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 80adec46b0972..cc0bb6a3c2cf0 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 1cf83aa651d7c..8efef8de29281 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -629,7 +629,8 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
return deinitImpl();
}
Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,
- StringRef InputTgtImage) {
+ StringRef InputTgtImage,
+ PluginContextTy *Context) {
ODBG(OLDT_Init) << "Load data from image "
<< static_cast<const void *>(InputTgtImage.bytes_begin());
@@ -657,7 +658,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;
@@ -1524,7 +1526,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_z...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/218387
More information about the llvm-commits
mailing list