[llvm] [offload][l0] Report actual driver version through OL_DEVICE_INFO_DRIVER_VERSION (PR #218412)
Jan Trusiłło via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 24 06:42:06 PDT 2026
https://github.com/311Volt created https://github.com/llvm/llvm-project/pull/218412
The Level Zero plugin currently reports the Level Zero API version as the driver version. Moreover, as the lookup table only includes API versions up to `1.12`, newer versions (e.g. `1.15`) end up being reported as Unknown.
Report the actual driver version instead:
- If the `zeIntelGetDriverVersionString` extension is available, use it. This string may be more detailed than what `zeDriverGetProperties` provides (e.g. hotfix suffix)
- If the above is not available, use `zeDriverGetProperties().driverVersion`, an U32 value whose octets make up major, minor & build numbers.
- The driver version is fetched once per context and stored in an `std::string` inside of the `L0ContextTy`.
As the PR uses L0 extensions, it uses the dispatch mechanism introduced in https://github.com/llvm/llvm-project/pull/218367 .
Stacked-on: https://github.com/llvm/llvm-project/pull/218367 . Review `9212312` and later.
Assisted-by: Codex (5.6 Sol)
>From 9083dfdd0cb8268f60ccb12699827e7a7a34c706 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Fri, 21 Aug 2026 16:22:41 +0200
Subject: [PATCH 1/4] [Offload] Implement better level zero dispatch
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 91 +----------------
.../level_zero/dynamic_l0/level_zero/ze_api.h | 2 +
.../level_zero/include/L0CmdListManager.h | 33 +++++--
.../level_zero/include/L0Compat.h | 13 +++
.../level_zero/include/L0Context.h | 99 ++++++++++++++++---
.../level_zero/src/L0Context.cpp | 57 ++++++++---
.../level_zero/src/L0Kernel.cpp | 4 +-
.../level_zero/src/L0Plugin.cpp | 4 +-
.../level_zero/src/L0Queue.cpp | 48 ++++++++-
9 files changed, 217 insertions(+), 134 deletions(-)
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) {
// Unlock KEnv lock after launching the kernel.
llvm::scope_exit UnlockGuard([&KEnv]() { KEnv.Lock.unlock(); });
- return CmdList->appendLaunchKernelWithArgs(
- Kernel, &KEnv.GroupCounts, &KEnv.GroupSizes, KEnv.ArgPtrs, SignalEvent,
- NumWaitEvents, WaitEvents, KEnv.IsCooperative);
+
+ bool AppendLaunchKernelWithArgsAvailable =
+ Device.getL0Context().LaunchKernelWithArguments.available();
+
+ if (AppendLaunchKernelWithArgsAvailable && Device.getL0Context().AppendLaunchKernelSupported.load(std::memory_order_acquire)) {
+ auto Result = CmdList->appendLaunchKernelWithArgs(
+ Kernel, &KEnv.GroupCounts, &KEnv.GroupSizes, KEnv.ArgPtrs, SignalEvent,
+ NumWaitEvents, WaitEvents, KEnv.IsCooperative);
+
+ // Can a context have multiple users?
+ if (!Device.getL0Context().AppendLaunchKernelSupported.load(std::memory_order_acquire)) {
+ // Commandlist failed to launch kernel with arguments, fallback to older
+ // API.
+ consumeError(std::move(Result));
+ } else {
+ return Result;
+ }
+ }
+
+ // Submit kernel using older set of APIs - zeKernelSetArgumentValue
+ auto &GroupSizes = KEnv.GroupSizes;
+ auto Res = zeKernelSetGroupSize(Kernel, GroupSizes.groupSizeX,
+ GroupSizes.groupSizeY, GroupSizes.groupSizeZ);
+ if (Res != ZE_RESULT_SUCCESS)
+ return error::createOffloadError(ErrorCode::UNKNOWN,
+ "Could not set group size!");
+
+ auto &KernelProperties = KEnv.KernelPR;
+
+ for (uint32_t KernelArg = 0; KernelArg < KernelProperties.NumKernelArgs;
+ KernelArg++) {
+ uint32_t ArgSize = KernelProperties.ArgSizes[KernelArg];
+
+ Res = zeKernelSetArgumentValue(Kernel, KernelArg, ArgSize,
+ KEnv.ArgPtrs[KernelArg]);
+ if (Res != ZE_RESULT_SUCCESS)
+ return error::createOffloadError(ErrorCode::UNKNOWN,
+ "Could not set argument to a kernel!");
+ }
+
+ return CmdList->appendLaunchKernel(Kernel, &KEnv.GroupCounts, SignalEvent,
+ NumWaitEvents, WaitEvents,
+ KEnv.IsCooperative);
}
Error L0QueueTy::memoryFill(void *Ptr, const void *Pattern, size_t PatternSize,
>From 919f6c931653ed8dae13465ca426cf74fc03ef5d Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Mon, 24 Aug 2026 12:23:35 +0200
Subject: [PATCH 2/4] Fix format
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 1 -
.../level_zero/include/L0CmdListManager.h | 3 ++-
offload/plugins-nextgen/level_zero/src/L0Queue.cpp | 9 ++++++---
3 files changed, 8 insertions(+), 5 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 0908051cd6d10..d09c936028a09 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -96,7 +96,6 @@ DLWRAP(zeCommandListAppendWaitOnEvents, 3)
DLWRAP(zeEventQueryStatus, 1)
DLWRAP(zeDriverGetDefaultContext, 1)
-
DLWRAP_FINALIZE()
#ifdef _WIN32
diff --git a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
index 70ec1726d3743..31c429b4bced9 100644
--- a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
+++ b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
@@ -148,7 +148,8 @@ class L0CmdListManagerTy {
WaitEvents);
if (Result == ze_result_t::ZE_RESULT_ERROR_UNSUPPORTED_FEATURE) {
- Context.AppendLaunchKernelSupported.store(false, std::memory_order_release);
+ Context.AppendLaunchKernelSupported.store(false,
+ std::memory_order_release);
return Plugin::error(
ErrorCode::UNSUPPORTED,
"zeCommandListAppendLaunchKernelWithArguments is not "
diff --git a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
index 5618f214428f6..18d57ac2cb748 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
@@ -57,14 +57,17 @@ Error L0QueueTy::dispatchLaunchKernel(ze_kernel_handle_t Kernel,
bool AppendLaunchKernelWithArgsAvailable =
Device.getL0Context().LaunchKernelWithArguments.available();
-
- if (AppendLaunchKernelWithArgsAvailable && Device.getL0Context().AppendLaunchKernelSupported.load(std::memory_order_acquire)) {
+
+ if (AppendLaunchKernelWithArgsAvailable &&
+ Device.getL0Context().AppendLaunchKernelSupported.load(
+ std::memory_order_acquire)) {
auto Result = CmdList->appendLaunchKernelWithArgs(
Kernel, &KEnv.GroupCounts, &KEnv.GroupSizes, KEnv.ArgPtrs, SignalEvent,
NumWaitEvents, WaitEvents, KEnv.IsCooperative);
// Can a context have multiple users?
- if (!Device.getL0Context().AppendLaunchKernelSupported.load(std::memory_order_acquire)) {
+ if (!Device.getL0Context().AppendLaunchKernelSupported.load(
+ std::memory_order_acquire)) {
// Commandlist failed to launch kernel with arguments, fallback to older
// API.
consumeError(std::move(Result));
>From a432ac4217d128018e64082f914f3823398c5366 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Mon, 24 Aug 2026 15:25:11 +0200
Subject: [PATCH 3/4] Rename `addFallbackFunction`
---
offload/plugins-nextgen/level_zero/include/L0Context.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/offload/plugins-nextgen/level_zero/include/L0Context.h b/offload/plugins-nextgen/level_zero/include/L0Context.h
index 2b554ac207dae..cdf96063455b7 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Context.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Context.h
@@ -95,7 +95,7 @@ template <auto Fn, auto UnsupportedValue> class ZeDispatcher {
return true;
}
- void addFallbackFunction(decltype(Fn) FallbackFunc) {
+ void setFallbackFunction(decltype(Fn) FallbackFunc) {
UsesFuncPtr = true;
FuncPtr = FallbackFunc;
}
>From 9212312ab85f3a3ea117075f86bbe5b50ef194c3 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jan=20Trusi=C5=82=C5=82o?= <jan.trusillo at intel.com>
Date: Mon, 24 Aug 2026 13:28:23 +0000
Subject: [PATCH 4/4] [offload][l0] report actual driver version
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 1 +
.../level_zero/dynamic_l0/level_zero/ze_api.h | 17 +++++-
.../level_zero/include/L0Compat.h | 4 ++
.../level_zero/include/L0Context.h | 11 ++++
.../level_zero/src/L0Context.cpp | 57 +++++++++++++++++++
.../level_zero/src/L0Device.cpp | 14 +----
6 files changed, 91 insertions(+), 13 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index d09c936028a09..70bc3aaf36512 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -57,6 +57,7 @@ DLWRAP(zeDeviceGetMemoryProperties, 3)
DLWRAP(zeDeviceGetCacheProperties, 3)
DLWRAP(zeDeviceGetGlobalTimestamps, 3)
DLWRAP(zeDriverGetApiVersion, 2)
+DLWRAP(zeDriverGetProperties, 2)
DLWRAP(zeDriverGetExtensionFunctionAddress, 3)
DLWRAP(zeDriverGetExtensionProperties, 3)
DLWRAP(zeEventCreate, 3)
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 5c9bcb9616247..be8ad9db758ee 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
@@ -334,11 +334,24 @@ typedef struct _ze_uuid_t {
uint8_t id[16];
} ze_uuid_t;
+/* Driver UUID size */
+#ifndef ZE_MAX_DRIVER_UUID_SIZE
+#define ZE_MAX_DRIVER_UUID_SIZE 16
+#endif
+
/* Driver UUID */
typedef struct _ze_driver_uuid_t {
- ze_uuid_t id;
+ uint8_t id[ZE_MAX_DRIVER_UUID_SIZE];
} ze_driver_uuid_t;
+/* Driver properties */
+typedef struct _ze_driver_properties_t {
+ ze_structure_type_t stype;
+ void *pNext;
+ ze_driver_uuid_t uuid;
+ uint32_t driverVersion;
+} ze_driver_properties_t;
+
/* Device UUID */
typedef struct _ze_device_uuid_t {
ze_uuid_t id;
@@ -666,6 +679,8 @@ ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGet(uint32_t *pCount,
ze_driver_handle_t *phDrivers);
ZE_APIEXPORT ze_result_t ZE_APICALL
zeDriverGetApiVersion(ze_driver_handle_t hDriver, ze_api_version_t *version);
+ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGetProperties(
+ ze_driver_handle_t hDriver, ze_driver_properties_t *pDriverProperties);
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(
diff --git a/offload/plugins-nextgen/level_zero/include/L0Compat.h b/offload/plugins-nextgen/level_zero/include/L0Compat.h
index a751b17abf0e3..c1580991bc5b4 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Compat.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Compat.h
@@ -38,4 +38,8 @@ API_HELPER_OPTIONAL(ze_result_t, zeCommandListAppendHostFunction,
API_HELPER_OPTIONAL(ze_context_handle_t, zeDriverGetDefaultContext,
ze_driver_handle_t hDriver)
+API_HELPER_OPTIONAL(ze_result_t, zeIntelGetDriverVersionString,
+ ze_driver_handle_t hDriver, char *pDriverVersion,
+ size_t *pVersionSize)
+
#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 cdf96063455b7..7aa77d2f34247 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Context.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Context.h
@@ -129,6 +129,11 @@ class L0ContextTy {
/// API version supported by the Level Zero driver.
ze_api_version_t APIVersion = ZE_API_VERSION_CURRENT;
+ /// Version of the Level Zero driver.
+ std::string DriverVersion;
+
+ Expected<std::string> tryGetIntelDriverVersion();
+
/// Imported external pointers. Track this only for user-directed
/// imports/releases.
llvm::DenseMap<uintptr_t, size_t> ImportedPtrs;
@@ -206,6 +211,9 @@ class L0ContextTy {
/// Return driver API version.
ze_api_version_t getDriverAPIVersion() const { return APIVersion; }
+ /// Return driver version.
+ const std::string &getDriverVersion() const { return DriverVersion; }
+
/// Return the event pool of this driver.
EventPoolTy &getEventPool() { return EventPool; }
const EventPoolTy &getEventPool() const { return EventPool; }
@@ -229,6 +237,9 @@ class L0ContextTy {
ZE_RESULT_ERROR_UNSUPPORTED_FEATURE>
CommandListAppendHostFunction;
ZeDispatcher<zeDriverGetDefaultContext, nullptr> DriverGetDefaultContext;
+ ZeDispatcher<zeIntelGetDriverVersionString,
+ ZE_RESULT_ERROR_UNSUPPORTED_FEATURE>
+ IntelGetDriverVersionString;
};
} // 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 08cd6b9f63324..6dfc0183ea2a2 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Context.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Context.cpp
@@ -13,14 +13,60 @@
#include "L0Context.h"
#include "L0Plugin.h"
+#include <cstring>
+
namespace llvm::omp::target::plugin {
+static ze_result_t ZE_APICALL getDriverVersionFromProperties(
+ ze_driver_handle_t zeDriver, char *DriverVersion, size_t *VersionSize) {
+ if (!VersionSize)
+ return ZE_RESULT_ERROR_INVALID_NULL_POINTER;
+
+ ze_driver_properties_t DriverProperties{};
+ DriverProperties.stype = ZE_STRUCTURE_TYPE_DRIVER_PROPERTIES;
+ ze_result_t Result = zeDriverGetProperties(zeDriver, &DriverProperties);
+ if (Result != ZE_RESULT_SUCCESS)
+ return Result;
+
+ uint32_t PackedVersion = DriverProperties.driverVersion;
+ std::string Version = std::to_string((PackedVersion & 0xFF000000) >> 24) +
+ "." +
+ std::to_string((PackedVersion & 0x00FF0000) >> 16) +
+ "." + std::to_string(PackedVersion & 0x0000FFFF);
+
+ if (!DriverVersion) {
+ *VersionSize = Version.size();
+ return ZE_RESULT_SUCCESS;
+ }
+
+ if (*VersionSize < Version.size()) {
+ *VersionSize = Version.size();
+ return ZE_RESULT_ERROR_INVALID_SIZE;
+ }
+
+ std::memcpy(DriverVersion, Version.data(), Version.size());
+ *VersionSize = Version.size();
+ return ZE_RESULT_SUCCESS;
+}
+
L0ContextTy::L0ContextTy(LevelZeroPluginTy &Plugin, ze_driver_handle_t zeDriver,
int32_t DriverId)
: Plugin(Plugin), zeDriver(zeDriver) {}
L0ContextTy::~L0ContextTy() = default;
+Expected<std::string> L0ContextTy::tryGetIntelDriverVersion() {
+ size_t VersionSize = 0;
+ CALL_ZE_RET_ERROR(IntelGetDriverVersionString, zeDriver, nullptr,
+ &VersionSize);
+ std::string Version(VersionSize, '\0');
+ CALL_ZE_RET_ERROR(IntelGetDriverVersionString, zeDriver, Version.data(),
+ &VersionSize);
+ if (!Version.empty() && Version.back() == '\0')
+ Version.pop_back();
+ return Version;
+}
+
Error L0ContextTy::init() {
auto CleanupOnError = [&]() {
if (zeContext) {
@@ -75,6 +121,11 @@ Error L0ContextTy::init() {
zeDriver, "zeCommandListAppendHostFunction");
DriverGetDefaultContext.tryLoadingExperimental(zeDriver,
"zeDriverGetDefaultContext");
+ if (!IntelGetDriverVersionString.tryLoadingExperimental(
+ zeDriver, "zeIntelGetDriverVersionString")) {
+ IntelGetDriverVersionString.setFallbackFunction(
+ getDriverVersionFromProperties);
+ }
ODBG(OLDT_Init) << "APIs supported by the context with added extensions: ";
ODBG(OLDT_Init) << " zeCommandListAppendLaunchKernelWithArguments: "
@@ -99,6 +150,12 @@ Error L0ContextTy::init() {
zeDriver, "zexCommandListAppendHostFunction");
}
+ auto DriverVersionOrErr = tryGetIntelDriverVersion();
+ if (!DriverVersionOrErr)
+ return DriverVersionOrErr.takeError();
+ DriverVersion = std::move(*DriverVersionOrErr);
+ ODBG(OLDT_Init) << "Driver version is " << DriverVersion;
+
DefaultUserCtx = std::make_unique<LevelZeroPluginContextTy>(
Plugin, /*Devices=*/llvm::ArrayRef<GenericDeviceTy *>{}, zeDriver,
zeContext, /*OwnsZeContext=*/false);
diff --git a/offload/plugins-nextgen/level_zero/src/L0Device.cpp b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
index 08c45b534e16c..74840ee6b1be1 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Device.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
@@ -428,12 +428,6 @@ const char *L0DeviceTy::getArchCStr() const {
}
}
-static const char *DriverVersionToStrTable[] = {
- "1.0", "1.1", "1.2", "1.3", "1.4", "1.5", "1.6",
- "1.7", "1.8", "1.9", "1.10", "1.11", "1.12"};
-constexpr size_t DriverVersionToStrTableSize =
- sizeof(DriverVersionToStrTable) / sizeof(DriverVersionToStrTable[0]);
-
Expected<InfoTreeNode> L0DeviceTy::obtainInfoImpl() {
InfoTreeNode Info;
Info.add("Device Number", getDeviceId());
@@ -442,12 +436,8 @@ Expected<InfoTreeNode> L0DeviceTy::obtainInfoImpl() {
Info.add("Device Type", "GPU", "", DeviceInfo::TYPE);
Info.add("Vendor", "Intel", "", DeviceInfo::VENDOR);
Info.add("Vendor ID", getVendorId(), "", DeviceInfo::VENDOR_ID);
- auto DriverVersion = getDriverAPIVersion();
- if (DriverVersion < DriverVersionToStrTableSize)
- Info.add("Driver Version", DriverVersionToStrTable[DriverVersion], "",
- DeviceInfo::DRIVER_VERSION);
- else
- Info.add("Driver Version", "Unknown", "", DeviceInfo::DRIVER_VERSION);
+ Info.add("Driver Version", L0Context.getDriverVersion(), "",
+ DeviceInfo::DRIVER_VERSION);
Info.add("Device PCI ID", getPCIId());
Info.add("Device UUID", getUuid().data());
Info.add("Number of total EUs", getNumEUs(), "",
More information about the llvm-commits
mailing list