[llvm] [Offload] Implement better level zero dispatch (PR #218367)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 24 03:21:54 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: blazej-smorawski
<details>
<summary>Changes</summary>
This PR implements better handling of level zero APIs which might act differently depending on level zero version:
Here's a small table that I based the PR on, you can see that some APIs are available in the loader level through direct linking or dlopen (D) or require call to `zeDriverGetExtensionFunctionAddress` (X)
```
25.18 25.22 25.35 25.44 25.48 26.22
zeDriverGetDefaultContext D X X X X D
zeCommandListAppendLaunchKernelWithArguments D X X X X D
zeCommandListAppendLaunchKernelWithParameters D D X X X D
zeCommandListAppendHostFunction . . . . X X
zexCommandListAppendHostFunction . . . X . .
```
Please note that all the possibly unavailable APIs are stored inside a context, because one platform might have multiple drivers with different features supported.
---
Patch is 22.24 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/218367.diff
9 Files Affected:
- (modified) offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp (+3-88)
- (modified) offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h (+2)
- (modified) offload/plugins-nextgen/level_zero/include/L0CmdListManager.h (+25-8)
- (modified) offload/plugins-nextgen/level_zero/include/L0Compat.h (+13)
- (modified) offload/plugins-nextgen/level_zero/include/L0Context.h (+84-15)
- (modified) offload/plugins-nextgen/level_zero/src/L0Context.cpp (+41-16)
- (modified) offload/plugins-nextgen/level_zero/src/L0Kernel.cpp (+2-2)
- (modified) offload/plugins-nextgen/level_zero/src/L0Plugin.cpp (+2-2)
- (modified) offload/plugins-nextgen/level_zero/src/L0Queue.cpp (+45-3)
``````````diff
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 55e4b6b856a96..0908051cd6d10 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -94,6 +94,8 @@ DLWRAP(zeCommandListHostSynchronize, 2)
DLWRAP(zeCommandListAppendSignalEvent, 2)
DLWRAP(zeCommandListAppendWaitOnEvents, 3)
DLWRAP(zeEventQueryStatus, 1)
+DLWRAP(zeDriverGetDefaultContext, 1)
+
DLWRAP_FINALIZE()
@@ -110,98 +112,11 @@ DLWRAP_FINALIZE()
#define DEBUG_PREFIX "TARGET " GETNAME(TARGET_NAME) " RTL"
#endif
-// Extension function pointer for getting argument sizes.
-static ze_result_t (*zexKernelGetArgumentSize_ptr)(ze_kernel_handle_t, uint32_t,
- uint32_t *) = nullptr;
-
-static bool zeCommandListAppendLaunchKernelWithArgumentsFallbackAvailable() {
- static std::once_flag zexKernelGetArgumentSize_once;
-
- // Load zexKernelGetArgumentSize extension if available.
- std::call_once(zexKernelGetArgumentSize_once, []() {
- uint32_t DriverCount = 0;
- if (zeDriverGet(&DriverCount, nullptr) == ZE_RESULT_SUCCESS &&
- DriverCount > 0) {
- ze_driver_handle_t Driver;
- DriverCount = 1;
- if (zeDriverGet(&DriverCount, &Driver) == ZE_RESULT_SUCCESS) {
- void *ExtFunc = nullptr;
- if (zeDriverGetExtensionFunctionAddress(
- Driver, "zexKernelGetArgumentSize", &ExtFunc) ==
- ZE_RESULT_SUCCESS &&
- ExtFunc) {
- zexKernelGetArgumentSize_ptr =
- reinterpret_cast<decltype(zexKernelGetArgumentSize_ptr)>(ExtFunc);
- ODBG(OLDT_Init) << "Loaded zexKernelGetArgumentSize extension";
- }
- }
- }
- });
- return zexKernelGetArgumentSize_ptr != nullptr;
-}
-
-static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
- ze_command_list_handle_t hCommandList, ze_kernel_handle_t hKernel,
- const ze_group_count_t groupCounts, const ze_group_size_t groupSizes,
- void **pArguments, const void *pNext, ze_event_handle_t hSignalEvent,
- uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) {
-
- ze_result_t Res;
- Res = zeKernelSetGroupSize(hKernel, groupSizes.groupSizeX,
- groupSizes.groupSizeY, groupSizes.groupSizeZ);
- if (Res != ZE_RESULT_SUCCESS)
- return Res;
-
- ze_kernel_properties_t KernelProps = {};
- KernelProps.stype = ZE_STRUCTURE_TYPE_KERNEL_PROPERTIES;
- Res = zeKernelGetProperties(hKernel, &KernelProps);
- if (Res != ZE_RESULT_SUCCESS)
- return Res;
-
- uint32_t NumKernelArgs = KernelProps.numKernelArgs;
-
- for (uint32_t KernelArg = 0; KernelArg < NumKernelArgs; KernelArg++) {
- uint32_t ArgSize = 0;
-
- Res = zexKernelGetArgumentSize_ptr(hKernel, KernelArg, &ArgSize);
- if (Res != ZE_RESULT_SUCCESS)
- return Res;
-
- Res = zeKernelSetArgumentValue(hKernel, KernelArg, ArgSize,
- pArguments[KernelArg]);
- if (Res != ZE_RESULT_SUCCESS)
- return Res;
- }
-
- bool IsCooperative = false;
- if (pNext) {
- const ze_command_list_append_launch_kernel_param_cooperative_desc_t
- *CoopDesc = static_cast<
- const ze_command_list_append_launch_kernel_param_cooperative_desc_t
- *>(pNext);
- if (CoopDesc->stype ==
- ZE_STRUCTURE_TYPE_COMMAND_LIST_APPEND_PARAM_COOPERATIVE_DESC)
- IsCooperative = CoopDesc->isCooperative;
- }
-
- if (IsCooperative)
- return zeCommandListAppendLaunchCooperativeKernel(
- hCommandList, hKernel, &groupCounts, hSignalEvent, numWaitEvents,
- phWaitEvents);
- return zeCommandListAppendLaunchKernel(hCommandList, hKernel, &groupCounts,
- hSignalEvent, numWaitEvents,
- phWaitEvents);
-}
-
static struct {
const char *Name;
void *FallbackFunc;
bool (*FallbackAvailable)();
-} ZeFallbacksTbl[] = {
- {"zeCommandListAppendLaunchKernelWithArguments",
- reinterpret_cast<void *>(
- &zeCommandListAppendLaunchKernelWithArgumentsFallback),
- zeCommandListAppendLaunchKernelWithArgumentsFallbackAvailable}};
+} ZeFallbacksTbl[] = {};
constexpr size_t ZeFallbacksTblSz =
sizeof(ZeFallbacksTbl) / sizeof(ZeFallbacksTbl[0]);
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h b/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
index 4d8cf1e9367c0..5c9bcb9616247 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
@@ -670,6 +670,8 @@ ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGetExtensionFunctionAddress(
ze_driver_handle_t hDriver, const char *name, void **ppFunctionAddress);
ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGetExtensionProperties(
ze_driver_handle_t hDriver, uint32_t *pCount, void *pExtensionProperties);
+ZE_APIEXPORT ze_context_handle_t ZE_APICALL
+zeDriverGetDefaultContext(ze_driver_handle_t hDriver);
/* Device functions */
ZE_APIEXPORT ze_result_t ZE_APICALL zeDeviceGet(ze_driver_handle_t hDriver,
diff --git a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
index 1f86006b30a0b..70ec1726d3743 100644
--- a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
+++ b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
@@ -18,6 +18,8 @@
#include "L0Defs.h"
#include "L0Trace.h"
#include "PluginInterface.h"
+#include "level_zero/ze_api.h"
+#include <atomic>
#include <mutex>
namespace llvm::omp::target::plugin {
@@ -129,7 +131,7 @@ class L0CmdListManagerTy {
ze_event_handle_t SignalEvent = nullptr, uint32_t NumWaitEvents = 0,
ze_event_handle_t *WaitEvents = nullptr, bool IsCooperative = false) {
- if (!api_helper::canCall<zeCommandListAppendLaunchKernelWithArguments>())
+ if (Context.LaunchKernelWithArguments.available() == false)
return Plugin::error(
ErrorCode::UNSUPPORTED,
"zeCommandListAppendLaunchKernelWithArguments is not "
@@ -140,10 +142,26 @@ class L0CmdListManagerTy {
static_cast<ze_bool_t>(IsCooperative)};
std::lock_guard<std::mutex> Lock(Mtx);
- CALL_ZE_RET_ERROR(zeCommandListAppendLaunchKernelWithArguments, CmdList,
- Kernel, *GroupCounts, *GroupSizes, ArgPtrs,
- IsCooperative ? &CoopDesc : nullptr, SignalEvent,
- NumWaitEvents, WaitEvents);
+ auto Result = Context.LaunchKernelWithArguments(
+ CmdList, Kernel, *GroupCounts, *GroupSizes, ArgPtrs,
+ IsCooperative ? &CoopDesc : nullptr, SignalEvent, NumWaitEvents,
+ WaitEvents);
+
+ if (Result == ze_result_t::ZE_RESULT_ERROR_UNSUPPORTED_FEATURE) {
+ Context.AppendLaunchKernelSupported.store(false, std::memory_order_release);
+ return Plugin::error(
+ ErrorCode::UNSUPPORTED,
+ "zeCommandListAppendLaunchKernelWithArguments is not "
+ "supported on this driver");
+ }
+
+ if (Result != ze_result_t::ZE_RESULT_SUCCESS) {
+ return Plugin::error(getOffloadErrorCode(Result),
+ "zeCommandListAppendLaunchKernelWithArguments "
+ "failed with error %d, %s",
+ Result, getZeErrorName(Result));
+ }
+
return Plugin::success();
}
@@ -177,13 +195,12 @@ class L0CmdListManagerTy {
ze_event_handle_t SignalEvent = nullptr,
uint32_t NumWaitEvents = 0,
ze_event_handle_t *WaitEvents = nullptr) {
- auto zeCommandListAppendHost = Context.zeCommandListAppendHostFunction;
- if (!zeCommandListAppendHost)
+ if (!Context.CommandListAppendHostFunction.available())
return Plugin::error(ErrorCode::UNSUPPORTED,
"zeCommandListAppendHostFunction extension is not "
"available on this driver");
std::lock_guard<std::mutex> Lock(Mtx);
- CALL_ZE_RET_ERROR(zeCommandListAppendHost, CmdList,
+ CALL_ZE_RET_ERROR(Context.CommandListAppendHostFunction, CmdList,
reinterpret_cast<void *>(Callback), UserData,
/*pReserved*/ nullptr, SignalEvent, NumWaitEvents,
WaitEvents);
diff --git a/offload/plugins-nextgen/level_zero/include/L0Compat.h b/offload/plugins-nextgen/level_zero/include/L0Compat.h
index dfa49e080033d..a751b17abf0e3 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Compat.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Compat.h
@@ -25,4 +25,17 @@ API_HELPER_OPTIONAL(ze_result_t, zeCommandListAppendLaunchKernelWithArguments,
const void *pNext, ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents)
+API_HELPER_OPTIONAL(ze_result_t, zexKernelGetArgumentSize,
+ ze_kernel_handle_t hKernel, uint32_t argIndex,
+ uint32_t *pArgSize)
+
+API_HELPER_OPTIONAL(ze_result_t, zeCommandListAppendHostFunction,
+ ze_command_list_handle_t hCommandList,
+ void *pfnHostFunction, void *pUserData, void *pReserved,
+ ze_event_handle_t hSignalEvent, uint32_t numWaitEvents,
+ ze_event_handle_t *phWaitEvents)
+
+API_HELPER_OPTIONAL(ze_context_handle_t, zeDriverGetDefaultContext,
+ ze_driver_handle_t hDriver)
+
#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0COMPAT_H
diff --git a/offload/plugins-nextgen/level_zero/include/L0Context.h b/offload/plugins-nextgen/level_zero/include/L0Context.h
index 09a19702baa23..2b554ac207dae 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Context.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Context.h
@@ -13,9 +13,12 @@
#ifndef OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0CONTEXT_H
#define OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0CONTEXT_H
+#include "APIHelpers.h"
+#include "L0Compat.h"
#include "L0Event.h"
#include "L0Memory.h"
#include "PerThreadTable.h"
+#include "level_zero/ze_api.h"
namespace llvm::omp::target::plugin {
@@ -32,6 +35,76 @@ class L0ContextTLSTy {
Error deinit() { return StagingBuffer.clear(); }
};
+// Helper for managing Level Zero APIs.
+// It provides two interfaces - by default it tries to call the function
+// directly - either through dlopen or directly linked (see L0DynWrapper.cpp).
+// It is also possible to call through an internal function pointer, which
+// can be populated using `tryLoadingExperimental` using
+// `zeDriverGetExtensionFunctionAddress` or simple set method
+// `addFallbackFunction`. It was implemented in order to support different
+// versions of level zero software stack and different kinds of drivers.
+template <auto Fn, auto UnsupportedValue> class ZeDispatcher {
+public:
+ constexpr ZeDispatcher() = default;
+
+ [[nodiscard]]
+ bool available() const {
+ if (UsesFuncPtr)
+ return FuncPtr != nullptr;
+
+ return api_helper::canCall<Fn>();
+ }
+
+ explicit operator bool() const { return available(); }
+
+ template <typename... Args>
+ decltype(auto) operator()(Args &&...ArgsList) const {
+ // Need to cast the type to avoid mismatch of return type deduction
+ using ReturnTy = std::invoke_result_t<decltype(Fn), Args...>;
+ if (UsesFuncPtr) {
+ if (FuncPtr == nullptr) {
+ return static_cast<ReturnTy>(UnsupportedValue);
+ }
+ auto Result = FuncPtr(std::forward<Args>(ArgsList)...);
+ return Result;
+ }
+
+ if (!api_helper::canCall<Fn>()) {
+ return static_cast<ReturnTy>(UnsupportedValue);
+ }
+ auto Result = Fn(std::forward<Args>(ArgsList)...);
+
+ return Result;
+ }
+
+ bool tryLoadingExperimental(ze_driver_handle_t zeDriver,
+ const char *FuncName) {
+ if (api_helper::canCall<Fn>()) {
+ return true; // Function is already available, no need to load it using
+ // experimental API.
+ }
+
+ auto Result = zeDriverGetExtensionFunctionAddress(
+ zeDriver, FuncName, reinterpret_cast<void **>(&FuncPtr));
+
+ if (Result != ZE_RESULT_SUCCESS || FuncPtr == nullptr) {
+ return false;
+ }
+
+ UsesFuncPtr = true;
+ return true;
+ }
+
+ void addFallbackFunction(decltype(Fn) FallbackFunc) {
+ UsesFuncPtr = true;
+ FuncPtr = FallbackFunc;
+ }
+
+private:
+ bool UsesFuncPtr = false;
+ decltype(Fn) FuncPtr = nullptr;
+};
+
struct L0ContextTLSTableTy
: public PerThreadContainer<
std::unordered_map<ze_context_handle_t, L0ContextTLSTy>> {
@@ -145,21 +218,17 @@ class L0ContextTy {
const MemAllocatorTy &getHostMemAllocator() const { return HostMemAllocator; }
MemAllocatorTy &getHostMemAllocator() { return HostMemAllocator; }
- /// Level Zero extension function pointer for kernel argument size query.
- ze_result_t(ZE_APICALL *zexKernelGetArgumentSize)(
- ze_kernel_handle_t hKernel, uint32_t argIndex,
- uint32_t *pArgSize) = nullptr;
-
- /// Level Zero extension function pointer for host function callbacks.
- ze_result_t(ZE_APICALL *zeCommandListAppendHostFunction)(
- ze_command_list_handle_t hCommandList, void *pfnHostFunction,
- void *pUserData, void *pReserved, ze_event_handle_t hSignalEvent,
- uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) = nullptr;
-
- /// Level Zero extension function pointer for querying the driver's default
- /// ze_context, when the extension is supported.
- ze_context_handle_t(ZE_APICALL *zeDriverGetDefaultContext)(
- ze_driver_handle_t hDriver) = nullptr;
+ std::atomic<bool> AppendLaunchKernelSupported = true;
+
+ ZeDispatcher<zeCommandListAppendLaunchKernelWithArguments,
+ ZE_RESULT_ERROR_UNSUPPORTED_FEATURE>
+ LaunchKernelWithArguments;
+ ZeDispatcher<zexKernelGetArgumentSize, ZE_RESULT_ERROR_UNSUPPORTED_FEATURE>
+ KernelGetArgumentSize;
+ ZeDispatcher<zeCommandListAppendHostFunction,
+ ZE_RESULT_ERROR_UNSUPPORTED_FEATURE>
+ CommandListAppendHostFunction;
+ ZeDispatcher<zeDriverGetDefaultContext, nullptr> DriverGetDefaultContext;
};
} // namespace llvm::omp::target::plugin
diff --git a/offload/plugins-nextgen/level_zero/src/L0Context.cpp b/offload/plugins-nextgen/level_zero/src/L0Context.cpp
index cdcd1210cbcf1..08cd6b9f63324 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Context.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Context.cpp
@@ -57,22 +57,47 @@ Error L0ContextTy::init() {
return Err;
}
- ze_result_t RC;
- CALL_ZE(RC, zeDriverGetExtensionFunctionAddress, zeDriver,
- "zexKernelGetArgumentSize", (void **)&zexKernelGetArgumentSize);
- if (RC != ZE_RESULT_SUCCESS)
- zexKernelGetArgumentSize = nullptr;
-
- CALL_ZE(RC, zeDriverGetExtensionFunctionAddress, zeDriver,
- "zeCommandListAppendHostFunction",
- (void **)&zeCommandListAppendHostFunction);
- if (RC != ZE_RESULT_SUCCESS)
- zeCommandListAppendHostFunction = nullptr;
-
- CALL_ZE(RC, zeDriverGetExtensionFunctionAddress, zeDriver,
- "zeDriverGetDefaultContext", (void **)&zeDriverGetDefaultContext);
- if (RC != ZE_RESULT_SUCCESS)
- zeDriverGetDefaultContext = nullptr;
+ ODBG(OLDT_Init) << "APIs supported by the context with dlopen: ";
+ ODBG(OLDT_Init) << " zeCommandListAppendLaunchKernelWithArguments: "
+ << (LaunchKernelWithArguments.available() ? "yes" : "no");
+ ODBG(OLDT_Init) << " zexKernelGetArgumentSize: "
+ << (KernelGetArgumentSize.available() ? "yes" : "no");
+ ODBG(OLDT_Init) << " zeCommandListAppendHostFunction: "
+ << (CommandListAppendHostFunction.available() ? "yes" : "no");
+ ODBG(OLDT_Init) << " zeDriverGetDefaultContext: "
+ << (DriverGetDefaultContext.available() ? "yes" : "no");
+
+ LaunchKernelWithArguments.tryLoadingExperimental(
+ zeDriver, "zeCommandListAppendLaunchKernelWithArguments");
+ KernelGetArgumentSize.tryLoadingExperimental(zeDriver,
+ "zexKernelGetArgumentSize");
+ CommandListAppendHostFunction.tryLoadingExperimental(
+ zeDriver, "zeCommandListAppendHostFunction");
+ DriverGetDefaultContext.tryLoadingExperimental(zeDriver,
+ "zeDriverGetDefaultContext");
+
+ ODBG(OLDT_Init) << "APIs supported by the context with added extensions: ";
+ ODBG(OLDT_Init) << " zeCommandListAppendLaunchKernelWithArguments: "
+ << (LaunchKernelWithArguments.available() ? "yes" : "no");
+ ODBG(OLDT_Init) << " zexKernelGetArgumentSize: "
+ << (KernelGetArgumentSize.available() ? "yes" : "no");
+ ODBG(OLDT_Init) << " zeCommandListAppendHostFunction: "
+ << (CommandListAppendHostFunction.available() ? "yes" : "no");
+ ODBG(OLDT_Init) << " zeDriverGetDefaultContext: "
+ << (DriverGetDefaultContext.available() ? "yes" : "no");
+
+ if (!LaunchKernelWithArguments.available() &&
+ KernelGetArgumentSize.available()) {
+ // Launch kernel was not available, both through dlopen and experimental API
+ // use fallback with KernelGetArgumentSize
+ // LaunchKernelWithArguments.addFallbackFunction(zeCommandListAppendLaunchKernelWithArgumentsFallback);
+ }
+
+ if (!CommandListAppendHostFunction.available()) {
+ // Try again with a name used in compute runtime 25.35 to 25.48
+ CommandListAppendHostFunction.tryLoadingExperimental(
+ zeDriver, "zexCommandListAppendHostFunction");
+ }
DefaultUserCtx = std::make_unique<LevelZeroPluginContextTy>(
Plugin, /*Devices=*/llvm::ArrayRef<GenericDeviceTy *>{}, zeDriver,
diff --git a/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp b/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp
index bd0e21d2301e9..69bcc40ee1f97 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp
@@ -43,10 +43,10 @@ Error L0KernelTy::readKernelProperties(L0ProgramTy &Program) {
// Query and cache argument sizes if extension is available.
auto &Context = L0Device.getL0Context();
- if (KernelPR.NumKernelArgs > 0 && Context.zexKernelGetArgumentSize) {
+ if (KernelPR.NumKernelArgs > 0 && Context.KernelGetArgumentSize.available()) {
KernelPR.ArgSizes = std::make_unique<uint32_t[]>(KernelPR.NumKernelArgs);
for (uint32_t I = 0; I < KernelPR.NumKernelArgs; I++) {
- CALL_ZE_RET_ERROR(Context.zexKernelGetArgumentSize, zeKernel, I,
+ CALL_ZE_RET_ERROR(Context.KernelGetArgumentSize, zeKernel, I,
&KernelPR.ArgSizes[I]);
}
}
diff --git a/offload/plugins-nextgen/level_zero/src/L0Plugin.cpp b/offload/plugins-nextgen/level_zero/src/L0Plugin.cpp
index a8c3e522c36ba..93c09bc7acee5 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Plugin.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Plugin.cpp
@@ -305,8 +305,8 @@ LevelZeroPluginTy::createPluginContext(
ze_context_handle_t ZeContext = nullptr;
bool OwnsZeContext = false;
- if (IsFullDriver && DriverCtx.zeDriverGetDefaultContext)
- ZeContext = DriverCtx.zeDriverGetDefaultContext(Driver);
+ if (IsFullDriver && DriverCtx.DriverGetDefaultContext.available())
+ ZeContext = DriverCtx.DriverGetDefaultContext(Driver);
if (!ZeContext) {
ze_context_desc_t Desc{ZE_STRUCTURE_TYPE_CONTEXT_DESC, nullptr, 0};
CALL_ZE_RET_ERROR(zeContextCreate, Driver, &Desc, &ZeContext);
diff --git a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
index 9382780c58238..5618f214428f6 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
@@ -14,7 +14,9 @@
#include "L0Device.h"
#include "L0Kernel.h"
#include "L0Plugin.h"
+#include "PluginInterface.h"
#include "llvm/ADT/ScopeExit.h"
+#include "llvm/Support/Error.h"
#include "llvm/Support/MathExtras.h"
#include <algorithm>
@@ -52,9 +54,49 @@ Error L0QueueTy::dispatchLaunchKernel(ze_kernel_handle_t Kernel,
ze_event_handle_t *WaitEvents)...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/218367
More information about the llvm-commits
mailing list