[llvm] [OFFLOAD][L0] Add support for dynamic l0 fallbacks (PR #200517)
Alex Duran via llvm-commits
llvm-commits at lists.llvm.org
Mon Jun 1 04:14:27 PDT 2026
https://github.com/adurang updated https://github.com/llvm/llvm-project/pull/200517
>From 1ea8a5257994326901873281555df83a4a4557d2 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Fri, 29 May 2026 15:40:05 -0700
Subject: [PATCH 1/4] [OFFLOAD][L0] Add support for dynamic l0 fallbacks
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 132 +++++++++++++++++-
1 file changed, 125 insertions(+), 7 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 67cbdca3ab82b..62a27588d0daf 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -106,6 +106,117 @@ 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 zexKernelGetArgumentSize_initialized = false;
+
+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;
+
+ if (!zexKernelGetArgumentSize_initialized) {
+ zexKernelGetArgumentSize_initialized = true;
+
+ // Get the driver to query for extensions
+ uint32_t driverCount = 0;
+ res = zeDriverGet(&driverCount, nullptr);
+ if (res == ZE_RESULT_SUCCESS && driverCount > 0) {
+ ze_driver_handle_t driver;
+ driverCount = 1;
+ res = zeDriverGet(&driverCount, &driver);
+ if (res == ZE_RESULT_SUCCESS) {
+ // Try to get the extension function address
+ void *extFunc = nullptr;
+ res = zeDriverGetExtensionFunctionAddress(
+ driver, "zexKernelGetArgumentSize", &extFunc);
+ if (res == ZE_RESULT_SUCCESS && extFunc) {
+ zexKernelGetArgumentSize_ptr =
+ reinterpret_cast<decltype(zexKernelGetArgumentSize_ptr)>(extFunc);
+ ODBG(OLDT_Init) << "Loaded zexKernelGetArgumentSize extension";
+ }
+ }
+ }
+ if (!zexKernelGetArgumentSize_ptr) {
+ ODBG(OLDT_Kernel)
+ << "zeCommandListAppendLaunchKernelWithArguments is not "
+ "available, and no fallback is possible without "
+ "argument size information.";
+ return ZE_RESULT_ERROR_UNSUPPORTED_FEATURE;
+ }
+ }
+
+ 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 i = 0; i < numKernelArgs; i++) {
+ uint32_t argSize = 0;
+
+ res = zexKernelGetArgumentSize_ptr(hKernel, i, &argSize);
+ if (res != ZE_RESULT_SUCCESS)
+ return res;
+
+ res = zeKernelSetArgumentValue(hKernel, i, argSize, pArguments[i]);
+ 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);
+ } else {
+ return zeCommandListAppendLaunchKernel(hCommandList, hKernel, &groupCounts,
+ hSignalEvent, numWaitEvents,
+ phWaitEvents);
+ }
+}
+
+static struct {
+ const char *name;
+ void *fallback_func;
+} zeFallbackFuncs[] = {
+ {"zeCommandListAppendLaunchKernelWithArguments",
+ reinterpret_cast<void *>(
+ &zeCommandListAppendLaunchKernelWithArgumentsFallback)}};
+constexpr size_t zeFallbackFuncsSz =
+ sizeof(zeFallbackFuncs) / sizeof(zeFallbackFuncs[0]);
+
+static void *findZeFallback(const char *name) {
+ for (size_t i = 0; i < zeFallbackFuncsSz; i++) {
+ if (strcmp(name, zeFallbackFuncs[i].name) == 0)
+ return zeFallbackFuncs[i].fallback_func;
+ }
+ return nullptr;
+}
+
static bool loadLevelZero() {
std::string L0Library{LEVEL_ZERO_LIBRARY};
std::string ErrMsg;
@@ -150,16 +261,23 @@ static bool loadLevelZero() {
const char *Sym = dlwrap::symbol(I);
void *P = DynlibHandle->getAddressOfSymbol(Sym);
+ void *Fallback = nullptr;
if (P == nullptr) {
- ODBG(OLDT_Init) << "Unable to find '" << Sym << "' in '" << L0Library
- << "'!";
- emitCheckVersion();
- return false;
+ Fallback = findZeFallback(Sym);
+ if (!Fallback) {
+ ODBG(OLDT_Init) << "Symbol '" << Sym << "' not found in '" << L0Library
+ << "' and no fallback is available!";
+ emitCheckVersion();
+ return false;
+ }
+ ODBG(OLDT_Init) << "Symbol '" << Sym << "' not found in '" << L0Library
+ << "'. Using fallback implementation -> " << Fallback;
}
- ODBG(OLDT_Init) << "Implementing " << Sym << " with dlsym(" << Sym
- << ") -> " << P;
+ if (P)
+ ODBG(OLDT_Init) << "Implementing " << Sym << " with dlsym(" << Sym
+ << ") -> " << P;
- *dlwrap::pointer(I) = P;
+ *dlwrap::pointer(I) = P ? P : Fallback;
}
return true;
>From 19274fc44fb970c28ca3d59af378e359aa6bae38 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Mon, 1 Jun 2026 04:03:43 -0700
Subject: [PATCH 2/4] fix naming conventions
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 76 +++++++++----------
1 file changed, 38 insertions(+), 38 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 62a27588d0daf..87f3025ae6131 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -106,7 +106,7 @@ DLWRAP_FINALIZE()
#define DEBUG_PREFIX "TARGET " GETNAME(TARGET_NAME) " RTL"
#endif
-// Extension function pointer for getting argument sizes
+// Extension function pointer for getting argument sizes.
static ze_result_t (*zexKernelGetArgumentSize_ptr)(ze_kernel_handle_t, uint32_t,
uint32_t *) = nullptr;
static bool zexKernelGetArgumentSize_initialized = false;
@@ -117,26 +117,26 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
void **pArguments, const void *pNext, ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) {
- ze_result_t res;
+ ze_result_t Res;
if (!zexKernelGetArgumentSize_initialized) {
zexKernelGetArgumentSize_initialized = true;
- // Get the driver to query for extensions
- uint32_t driverCount = 0;
- res = zeDriverGet(&driverCount, nullptr);
- if (res == ZE_RESULT_SUCCESS && driverCount > 0) {
- ze_driver_handle_t driver;
- driverCount = 1;
- res = zeDriverGet(&driverCount, &driver);
- if (res == ZE_RESULT_SUCCESS) {
+ // Get the driver to query for extensions.
+ uint32_t DriverCount = 0;
+ Res = zeDriverGet(&DriverCount, nullptr);
+ if (Res == ZE_RESULT_SUCCESS && DriverCount > 0) {
+ ze_driver_handle_t Driver;
+ DriverCount = 1;
+ Res = zeDriverGet(&DriverCount, &Driver);
+ if (Res == ZE_RESULT_SUCCESS) {
// Try to get the extension function address
- void *extFunc = nullptr;
- res = zeDriverGetExtensionFunctionAddress(
- driver, "zexKernelGetArgumentSize", &extFunc);
- if (res == ZE_RESULT_SUCCESS && extFunc) {
+ void *ExtFunc = nullptr;
+ Res = zeDriverGetExtensionFunctionAddress(
+ Driver, "zexKernelGetArgumentSize", &ExtFunc);
+ if (Res == ZE_RESULT_SUCCESS && ExtFunc) {
zexKernelGetArgumentSize_ptr =
- reinterpret_cast<decltype(zexKernelGetArgumentSize_ptr)>(extFunc);
+ reinterpret_cast<decltype(zexKernelGetArgumentSize_ptr)>(ExtFunc);
ODBG(OLDT_Init) << "Loaded zexKernelGetArgumentSize extension";
}
}
@@ -150,33 +150,33 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
}
}
- res = zeKernelSetGroupSize(hKernel, groupSizes.groupSizeX,
+ Res = zeKernelSetGroupSize(hKernel, groupSizes.groupSizeX,
groupSizes.groupSizeY, groupSizes.groupSizeZ);
- if (res != ZE_RESULT_SUCCESS)
- return res;
+ 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;
+ Res = zeKernelGetProperties(hKernel, &kernelProps);
+ if (Res != ZE_RESULT_SUCCESS)
+ return Res;
- uint32_t numKernelArgs = kernelProps.numKernelArgs;
+ uint32_t NumKernelArgs = kernelProps.numKernelArgs;
- for (uint32_t i = 0; i < numKernelArgs; i++) {
+ for (uint32_t i = 0; i < NumKernelArgs; i++) {
uint32_t argSize = 0;
- res = zexKernelGetArgumentSize_ptr(hKernel, i, &argSize);
- if (res != ZE_RESULT_SUCCESS)
- return res;
+ Res = zexKernelGetArgumentSize_ptr(hKernel, i, &argSize);
+ if (Res != ZE_RESULT_SUCCESS)
+ return Res;
- res = zeKernelSetArgumentValue(hKernel, i, argSize, pArguments[i]);
- if (res != ZE_RESULT_SUCCESS) {
- return res;
+ Res = zeKernelSetArgumentValue(hKernel, i, argSize, pArguments[i]);
+ if (Res != ZE_RESULT_SUCCESS) {
+ return Res;
}
}
- bool isCooperative = false;
+ bool IsCooperative = false;
if (pNext) {
const ze_command_list_append_launch_kernel_param_cooperative_desc_t
*coopDesc = static_cast<
@@ -184,11 +184,11 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
*>(pNext);
if (coopDesc->stype ==
ZE_STRUCTURE_TYPE_COMMAND_LIST_APPEND_PARAM_COOPERATIVE_DESC) {
- isCooperative = coopDesc->isCooperative;
+ IsCooperative = coopDesc->isCooperative;
}
}
- if (isCooperative) {
+ if (IsCooperative) {
return zeCommandListAppendLaunchCooperativeKernel(
hCommandList, hKernel, &groupCounts, hSignalEvent, numWaitEvents,
phWaitEvents);
@@ -202,17 +202,17 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
static struct {
const char *name;
void *fallback_func;
-} zeFallbackFuncs[] = {
+} ZeFallbacksTbl[] = {
{"zeCommandListAppendLaunchKernelWithArguments",
reinterpret_cast<void *>(
&zeCommandListAppendLaunchKernelWithArgumentsFallback)}};
-constexpr size_t zeFallbackFuncsSz =
- sizeof(zeFallbackFuncs) / sizeof(zeFallbackFuncs[0]);
+constexpr size_t ZeFallbacksTblSz =
+ sizeof(ZeFallbacksTbl) / sizeof(ZeFallbacksTbl[0]);
static void *findZeFallback(const char *name) {
- for (size_t i = 0; i < zeFallbackFuncsSz; i++) {
- if (strcmp(name, zeFallbackFuncs[i].name) == 0)
- return zeFallbackFuncs[i].fallback_func;
+ for (size_t i = 0; i < ZeFallbacksTblSz; i++) {
+ if (strcmp(name, ZeFallbacksTbl[i].name) == 0)
+ return ZeFallbacksTbl[i].fallback_func;
}
return nullptr;
}
>From 6d2e592047cced091f7a18ac6ac3b1f3808655e5 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Mon, 1 Jun 2026 04:09:13 -0700
Subject: [PATCH 3/4] make zexKernelGetArgumentSize_ptr init thread-safe
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 38 +++++++++----------
1 file changed, 18 insertions(+), 20 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 87f3025ae6131..bb09d05aca71b 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -7,6 +7,7 @@
#include <level_zero/ze_api.h>
#include <level_zero/zes_api.h>
#include <memory>
+#include <mutex>
#include "DLWrap.h"
#include "Shared/Debug.h"
@@ -109,7 +110,6 @@ DLWRAP_FINALIZE()
// Extension function pointer for getting argument sizes.
static ze_result_t (*zexKernelGetArgumentSize_ptr)(ze_kernel_handle_t, uint32_t,
uint32_t *) = nullptr;
-static bool zexKernelGetArgumentSize_initialized = false;
static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
ze_command_list_handle_t hCommandList, ze_kernel_handle_t hKernel,
@@ -117,37 +117,35 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
void **pArguments, const void *pNext, ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) {
+ static std::once_flag zexKernelGetArgumentSize_once;
ze_result_t Res;
- if (!zexKernelGetArgumentSize_initialized) {
- zexKernelGetArgumentSize_initialized = true;
-
- // Get the driver to query for extensions.
+ // Load zexKernelGetArgumentSize extension if available.
+ std::call_once(zexKernelGetArgumentSize_once, []() {
uint32_t DriverCount = 0;
- Res = zeDriverGet(&DriverCount, nullptr);
- if (Res == ZE_RESULT_SUCCESS && DriverCount > 0) {
+ if (zeDriverGet(&DriverCount, nullptr) == ZE_RESULT_SUCCESS &&
+ DriverCount > 0) {
ze_driver_handle_t Driver;
DriverCount = 1;
- Res = zeDriverGet(&DriverCount, &Driver);
- if (Res == ZE_RESULT_SUCCESS) {
- // Try to get the extension function address
+ if (zeDriverGet(&DriverCount, &Driver) == ZE_RESULT_SUCCESS) {
void *ExtFunc = nullptr;
- Res = zeDriverGetExtensionFunctionAddress(
- Driver, "zexKernelGetArgumentSize", &ExtFunc);
- if (Res == ZE_RESULT_SUCCESS && ExtFunc) {
+ if (zeDriverGetExtensionFunctionAddress(
+ Driver, "zexKernelGetArgumentSize", &ExtFunc) ==
+ ZE_RESULT_SUCCESS &&
+ ExtFunc) {
zexKernelGetArgumentSize_ptr =
reinterpret_cast<decltype(zexKernelGetArgumentSize_ptr)>(ExtFunc);
ODBG(OLDT_Init) << "Loaded zexKernelGetArgumentSize extension";
}
}
}
- if (!zexKernelGetArgumentSize_ptr) {
- ODBG(OLDT_Kernel)
- << "zeCommandListAppendLaunchKernelWithArguments is not "
- "available, and no fallback is possible without "
- "argument size information.";
- return ZE_RESULT_ERROR_UNSUPPORTED_FEATURE;
- }
+ });
+ if (!zexKernelGetArgumentSize_ptr) {
+ ODBG(OLDT_Kernel)
+ << "zeCommandListAppendLaunchKernelWithArguments is not "
+ "available, and no fallback is possible without "
+ "argument size information.";
+ return ZE_RESULT_ERROR_UNSUPPORTED_FEATURE;
}
Res = zeKernelSetGroupSize(hKernel, groupSizes.groupSizeX,
>From 9425d7d07b73aa4ef34449b45bded72f5b7eb0b3 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Mon, 1 Jun 2026 04:14:14 -0700
Subject: [PATCH 4/4] format
---
.../plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp | 7 +++----
1 file changed, 3 insertions(+), 4 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index bb09d05aca71b..8760d1acc3abc 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -141,10 +141,9 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
}
});
if (!zexKernelGetArgumentSize_ptr) {
- ODBG(OLDT_Kernel)
- << "zeCommandListAppendLaunchKernelWithArguments is not "
- "available, and no fallback is possible without "
- "argument size information.";
+ ODBG(OLDT_Kernel) << "zeCommandListAppendLaunchKernelWithArguments is not "
+ "available, and no fallback is possible without "
+ "argument size information.";
return ZE_RESULT_ERROR_UNSUPPORTED_FEATURE;
}
More information about the llvm-commits
mailing list