[llvm] [Offload] Implement better level zero dispatch (PR #218367)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Aug 26 02:42:51 PDT 2026
https://github.com/blazej-smorawski updated https://github.com/llvm/llvm-project/pull/218367
>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 cd200b539188298895a56d1e08cdeb17d500b024 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 26 Aug 2026 11:41:36 +0200
Subject: [PATCH 4/4] Remove fallback from dispatcher
---
offload/plugins-nextgen/level_zero/include/L0Context.h | 7 ++-----
offload/plugins-nextgen/level_zero/src/L0Context.cpp | 7 -------
2 files changed, 2 insertions(+), 12 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/include/L0Context.h b/offload/plugins-nextgen/level_zero/include/L0Context.h
index cdf96063455b7..5f8913625c86d 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Context.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Context.h
@@ -19,6 +19,8 @@
#include "L0Memory.h"
#include "PerThreadTable.h"
#include "level_zero/ze_api.h"
+#include <cassert>
+#include <cstdint>
namespace llvm::omp::target::plugin {
@@ -95,11 +97,6 @@ template <auto Fn, auto UnsupportedValue> class ZeDispatcher {
return true;
}
- void setFallbackFunction(decltype(Fn) FallbackFunc) {
- UsesFuncPtr = true;
- FuncPtr = FallbackFunc;
- }
-
private:
bool UsesFuncPtr = false;
decltype(Fn) FuncPtr = nullptr;
diff --git a/offload/plugins-nextgen/level_zero/src/L0Context.cpp b/offload/plugins-nextgen/level_zero/src/L0Context.cpp
index 08cd6b9f63324..38f7bd0fd0bad 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Context.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Context.cpp
@@ -86,13 +86,6 @@ Error L0ContextTy::init() {
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(
More information about the llvm-commits
mailing list