[llvm] [offload] Add `dlwrap::loaded` function to check for optional symbols (PR #210737)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 3 02:47:45 PDT 2026
https://github.com/blazej-smorawski updated https://github.com/llvm/llvm-project/pull/210737
>From 22ef2ce77a38d5b299d541af8c11dd545d856a77 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Fri, 3 Jul 2026 13:48:27 +0200
Subject: [PATCH 01/20] [offload] Add `dlwrap::loaded` function to check for
optional symbols
---
offload/plugins-nextgen/common/include/DLWrap.h | 5 +++++
1 file changed, 5 insertions(+)
diff --git a/offload/plugins-nextgen/common/include/DLWrap.h b/offload/plugins-nextgen/common/include/DLWrap.h
index 95ce86e123cd3..116cd4afe37a6 100644
--- a/offload/plugins-nextgen/common/include/DLWrap.h
+++ b/offload/plugins-nextgen/common/include/DLWrap.h
@@ -135,6 +135,8 @@ template <size_t Requested, size_t Required> constexpr void verboseAssert() {
static_assert(Requested == Required, "Arity Error");
}
+template <auto Fn> bool loaded();
+
} // namespace dlwrap
#define DLWRAP_INSTANTIATE(SYM_DEF, SYM_USE, ARITY) \
@@ -167,6 +169,9 @@ template <size_t Requested, size_t Required> constexpr void verboseAssert() {
return reinterpret_cast<T::FunctionType>(P); \
} \
}; \
+ template <> bool loaded<&(::SYMBOL)>() { \
+ return SYMBOL##_Trait::get() != nullptr; \
+ } \
}
#define DLWRAP_IMPL(SYMBOL, ARITY) \
>From 2b54c4597d0998f42aa9472f964ef205cd162fc7 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 22 Jul 2026 12:53:47 +0200
Subject: [PATCH 02/20] [offload] Add helpers to enable optional symbols from
external APIs
---
.../common/include/APIHelpers.h | 41 +++++++++++++++++++
.../plugins-nextgen/common/include/DLWrap.h | 12 +++++-
.../level_zero/include/L0CmdListManager.h | 10 +++++
.../level_zero/include/L0Compat.h | 28 +++++++++++++
4 files changed, 89 insertions(+), 2 deletions(-)
create mode 100644 offload/plugins-nextgen/common/include/APIHelpers.h
create mode 100644 offload/plugins-nextgen/level_zero/include/L0Compat.h
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
new file mode 100644
index 0000000000000..7c0e283ab5805
--- /dev/null
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -0,0 +1,41 @@
+//===-- Shared/APIHelpers.h - helpers for external APIs --*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// The header contains helper functions to make interactions with external APIs
+// such as CUDA or level zero easier
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_APIHELPERS_H
+#define OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_APIHELPERS_H
+
+#include "DLWrap.h"
+
+// Macro to mark external symbol as weak, so linker will be okay
+// if the symbol is missing. For direct linking (dlwrap::IsDlOpened<&name> ==
+// false), we need to check if linker could find the symbol. For symbols loaded
+// using dlsym we use dlwrap::loaded<name>().
+#define API_HELPER_OPTIONAL(return_type, name, ...) \
+ extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
+ template <> inline bool api_helper::canCall<name>() { \
+ if constexpr (dlwrap::IsDlOpened<&name>) \
+ return dlwrap::loaded<name>(); \
+ return name != nullptr; \
+ }
+
+namespace api_helper {
+
+// Default template specialization for extra safety
+template <auto Fn>
+bool canCall() {
+ static_assert(false, "api_helper::canCall() should only be called on symbols decorated with API_HELPER_OPTIONAL!");
+}
+
+} // namespace api_helper
+
+#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_APIHELPERS_H
\ No newline at end of file
diff --git a/offload/plugins-nextgen/common/include/DLWrap.h b/offload/plugins-nextgen/common/include/DLWrap.h
index 116cd4afe37a6..863d8006f246a 100644
--- a/offload/plugins-nextgen/common/include/DLWrap.h
+++ b/offload/plugins-nextgen/common/include/DLWrap.h
@@ -78,7 +78,7 @@
static size_t size(); \
static const char *symbol(size_t); /* get symbol name in [0, size()) */ \
static void ** \
- pointer(size_t); /* get pointer to function pointer in [0, size()) */ \
+ pointer(size_t); /* get pointer to function pointer in [0, size()) */ \
}
// DLWRAP_FINALIZE() implements the functions from DLWRAP_INITIALIZE
@@ -107,7 +107,9 @@ template <size_t S> struct count {
static constexpr size_t N = count<S - 1>::N;
};
-template <> struct count<0> { static constexpr size_t N = 0; };
+template <> struct count<0> {
+ static constexpr size_t N = 0;
+};
// Get a constexpr size_t ID, starts at zero
#define DLWRAP_ID() (dlwrap::type::count<__LINE__>::N)
@@ -135,8 +137,13 @@ template <size_t Requested, size_t Required> constexpr void verboseAssert() {
static_assert(Requested == Required, "Arity Error");
}
+// Template to check if a symbol was loaded successfully
template <auto Fn> bool loaded();
+// Template to check if a symbol is provided by dlwrap
+// By default all symbols resolve to false
+template <auto Fn> constexpr bool IsDlOpened = false;
+
} // namespace dlwrap
#define DLWRAP_INSTANTIATE(SYM_DEF, SYM_USE, ARITY) \
@@ -160,6 +167,7 @@ template <auto Fn> bool loaded();
DLWRAP_INC(); \
DLWRAP_SYMBOL(SYMBOL, DLWRAP_ID() - 1); \
namespace dlwrap { \
+ template <> inline constexpr bool IsDlOpened<&::SYMBOL> = true; \
struct SYMBOL##_Trait : public dlwrap::trait<decltype(&SYMBOL)> { \
using T = dlwrap::trait<decltype(&SYMBOL)>; \
static T::FunctionType get() { \
diff --git a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
index 94ddef231ff11..1f86006b30a0b 100644
--- a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
+++ b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
@@ -13,9 +13,11 @@
#ifndef OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0CMDLISTMANAGER_H
#define OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0CMDLISTMANAGER_H
+#include "L0Compat.h"
#include "L0Context.h"
#include "L0Defs.h"
#include "L0Trace.h"
+#include "PluginInterface.h"
#include <mutex>
namespace llvm::omp::target::plugin {
@@ -126,10 +128,18 @@ class L0CmdListManagerTy {
const ze_group_size_t *GroupSizes, void **ArgPtrs,
ze_event_handle_t SignalEvent = nullptr, uint32_t NumWaitEvents = 0,
ze_event_handle_t *WaitEvents = nullptr, bool IsCooperative = false) {
+
+ if (!api_helper::canCall<zeCommandListAppendLaunchKernelWithArguments>())
+ return Plugin::error(
+ ErrorCode::UNSUPPORTED,
+ "zeCommandListAppendLaunchKernelWithArguments is not "
+ "available on this driver");
+
ze_command_list_append_launch_kernel_param_cooperative_desc_t CoopDesc = {
ZE_STRUCTURE_TYPE_COMMAND_LIST_APPEND_PARAM_COOPERATIVE_DESC, nullptr,
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,
diff --git a/offload/plugins-nextgen/level_zero/include/L0Compat.h b/offload/plugins-nextgen/level_zero/include/L0Compat.h
new file mode 100644
index 0000000000000..2062c9624ad20
--- /dev/null
+++ b/offload/plugins-nextgen/level_zero/include/L0Compat.h
@@ -0,0 +1,28 @@
+//===--- Level Zero Target RTL Implementation -----------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Level Zero compatibility layer enabling us to compile using new APIs.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0COMPAT_H
+#define OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0COMPAT_H
+
+#include "APIHelpers.h"
+
+#include <level_zero/ze_api.h>
+
+API_HELPER_OPTIONAL(ze_result_t, zeCommandListAppendLaunchKernelWithArguments,
+ 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);
+
+#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0COMPAT_H
>From d08df4828fbb05a5e3e6f6375d127f1ad279b146 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 22 Jul 2026 14:43:45 +0200
Subject: [PATCH 03/20] [offload] Make `canCall` work on older compilers
---
offload/plugins-nextgen/common/include/APIHelpers.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 7c0e283ab5805..2bf331507ec8d 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -33,7 +33,7 @@ namespace api_helper {
// Default template specialization for extra safety
template <auto Fn>
bool canCall() {
- static_assert(false, "api_helper::canCall() should only be called on symbols decorated with API_HELPER_OPTIONAL!");
+ static_assert(sizeof(decltype(Fn)*) == 0, "api_helper::canCall() should only be called on symbols decorated with API_HELPER_OPTIONAL!");
}
} // namespace api_helper
>From fe88bd1418c9516671335b729cb02a34d421c889 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 22 Jul 2026 14:50:00 +0200
Subject: [PATCH 04/20] [offload] Fix format
---
offload/plugins-nextgen/common/include/APIHelpers.h | 9 +++++----
offload/plugins-nextgen/common/include/DLWrap.h | 2 +-
2 files changed, 6 insertions(+), 5 deletions(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 2bf331507ec8d..88a430bdb4839 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -22,7 +22,7 @@
// using dlsym we use dlwrap::loaded<name>().
#define API_HELPER_OPTIONAL(return_type, name, ...) \
extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
- template <> inline bool api_helper::canCall<name>() { \
+ template <> inline bool api_helper::canCall<name>() { \
if constexpr (dlwrap::IsDlOpened<&name>) \
return dlwrap::loaded<name>(); \
return name != nullptr; \
@@ -31,9 +31,10 @@
namespace api_helper {
// Default template specialization for extra safety
-template <auto Fn>
-bool canCall() {
- static_assert(sizeof(decltype(Fn)*) == 0, "api_helper::canCall() should only be called on symbols decorated with API_HELPER_OPTIONAL!");
+template <auto Fn> bool canCall() {
+ static_assert(sizeof(decltype(Fn) *) == 0,
+ "api_helper::canCall() should only be called on symbols "
+ "decorated with API_HELPER_OPTIONAL!");
}
} // namespace api_helper
diff --git a/offload/plugins-nextgen/common/include/DLWrap.h b/offload/plugins-nextgen/common/include/DLWrap.h
index 863d8006f246a..160731a176d6d 100644
--- a/offload/plugins-nextgen/common/include/DLWrap.h
+++ b/offload/plugins-nextgen/common/include/DLWrap.h
@@ -78,7 +78,7 @@
static size_t size(); \
static const char *symbol(size_t); /* get symbol name in [0, size()) */ \
static void ** \
- pointer(size_t); /* get pointer to function pointer in [0, size()) */ \
+ pointer(size_t); /* get pointer to function pointer in [0, size()) */ \
}
// DLWRAP_FINALIZE() implements the functions from DLWRAP_INITIALIZE
>From e5e08dac4867e3b7bf4bafeb9ea1aa8bd1190510 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Thu, 23 Jul 2026 16:16:08 +0200
Subject: [PATCH 05/20] [offload] Fix incorrect constexpr in
`API_HELPER_OPTIONAL`
---
offload/plugins-nextgen/common/include/APIHelpers.h | 11 +++++++----
offload/plugins-nextgen/common/include/DLWrap.h | 12 +++++-------
2 files changed, 12 insertions(+), 11 deletions(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 88a430bdb4839..85bdfb8f640e1 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -20,12 +20,15 @@
// if the symbol is missing. For direct linking (dlwrap::IsDlOpened<&name> ==
// false), we need to check if linker could find the symbol. For symbols loaded
// using dlsym we use dlwrap::loaded<name>().
-#define API_HELPER_OPTIONAL(return_type, name, ...) \
+#define API_HELPER_OPTIONAL(return_type, name, ...) \
extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
template <> inline bool api_helper::canCall<name>() { \
- if constexpr (dlwrap::IsDlOpened<&name>) \
- return dlwrap::loaded<name>(); \
- return name != nullptr; \
+ if (name == nullptr) \
+ /* Not loaded weak symbol */ \
+ return false; \
+ /* Symbols from dlwrap are never nullptr, but `loaded` might return false \
+ */ \
+ return dlwrap::loaded<name>(); \
}
namespace api_helper {
diff --git a/offload/plugins-nextgen/common/include/DLWrap.h b/offload/plugins-nextgen/common/include/DLWrap.h
index 160731a176d6d..394891abded9d 100644
--- a/offload/plugins-nextgen/common/include/DLWrap.h
+++ b/offload/plugins-nextgen/common/include/DLWrap.h
@@ -137,12 +137,11 @@ template <size_t Requested, size_t Required> constexpr void verboseAssert() {
static_assert(Requested == Required, "Arity Error");
}
-// Template to check if a symbol was loaded successfully
-template <auto Fn> bool loaded();
-
-// Template to check if a symbol is provided by dlwrap
-// By default all symbols resolve to false
-template <auto Fn> constexpr bool IsDlOpened = false;
+// Template to check if a symbol was loaded successfully.
+// Returns true for symbols that were not wrapped by dlwrap.
+template <auto Fn> bool loaded() {
+ return true;
+}
} // namespace dlwrap
@@ -167,7 +166,6 @@ template <auto Fn> constexpr bool IsDlOpened = false;
DLWRAP_INC(); \
DLWRAP_SYMBOL(SYMBOL, DLWRAP_ID() - 1); \
namespace dlwrap { \
- template <> inline constexpr bool IsDlOpened<&::SYMBOL> = true; \
struct SYMBOL##_Trait : public dlwrap::trait<decltype(&SYMBOL)> { \
using T = dlwrap::trait<decltype(&SYMBOL)>; \
static T::FunctionType get() { \
>From 73e60ffdfa92fb7dc5909464a102a7a1a8a7a7c3 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Thu, 23 Jul 2026 16:22:08 +0200
Subject: [PATCH 06/20] [offload] Fix format
---
offload/plugins-nextgen/common/include/APIHelpers.h | 4 ++--
offload/plugins-nextgen/common/include/DLWrap.h | 4 +---
2 files changed, 3 insertions(+), 5 deletions(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 85bdfb8f640e1..182f11fc4bc3a 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -20,12 +20,12 @@
// if the symbol is missing. For direct linking (dlwrap::IsDlOpened<&name> ==
// false), we need to check if linker could find the symbol. For symbols loaded
// using dlsym we use dlwrap::loaded<name>().
-#define API_HELPER_OPTIONAL(return_type, name, ...) \
+#define API_HELPER_OPTIONAL(return_type, name, ...) \
extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
template <> inline bool api_helper::canCall<name>() { \
if (name == nullptr) \
/* Not loaded weak symbol */ \
- return false; \
+ return false; \
/* Symbols from dlwrap are never nullptr, but `loaded` might return false \
*/ \
return dlwrap::loaded<name>(); \
diff --git a/offload/plugins-nextgen/common/include/DLWrap.h b/offload/plugins-nextgen/common/include/DLWrap.h
index 394891abded9d..16f5a86911da3 100644
--- a/offload/plugins-nextgen/common/include/DLWrap.h
+++ b/offload/plugins-nextgen/common/include/DLWrap.h
@@ -139,9 +139,7 @@ template <size_t Requested, size_t Required> constexpr void verboseAssert() {
// Template to check if a symbol was loaded successfully.
// Returns true for symbols that were not wrapped by dlwrap.
-template <auto Fn> bool loaded() {
- return true;
-}
+template <auto Fn> bool loaded() { return true; }
} // namespace dlwrap
>From 3127fbd29ef9f163872e04a5fb6a43612910f692 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Fri, 24 Jul 2026 10:34:31 +0200
Subject: [PATCH 07/20] [offload] Disable direct linking on Windows
---
offload/plugins-nextgen/amdgpu/CMakeLists.txt | 8 +++++++-
offload/plugins-nextgen/cuda/CMakeLists.txt | 8 +++++++-
offload/plugins-nextgen/level_zero/CMakeLists.txt | 8 +++++++-
3 files changed, 21 insertions(+), 3 deletions(-)
diff --git a/offload/plugins-nextgen/amdgpu/CMakeLists.txt b/offload/plugins-nextgen/amdgpu/CMakeLists.txt
index af2db85f74c64..4f680087e3fcf 100644
--- a/offload/plugins-nextgen/amdgpu/CMakeLists.txt
+++ b/offload/plugins-nextgen/amdgpu/CMakeLists.txt
@@ -8,7 +8,13 @@ target_sources(omptarget.rtl.amdgpu PRIVATE src/rtl.cpp)
target_include_directories(omptarget.rtl.amdgpu PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/utils)
-if(hsa-runtime64_FOUND AND NOT "amdgpu" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS)
+if (WIN32 AND NOT "amdgpu" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS)
+ message(WARNING "Direct linking against libhsa is not supported on Windows; "
+ "falling back to dlopen. Add 'amdgpu' to "
+ "LIBOMPTARGET_DLOPEN_PLUGINS to silence this warning.")
+endif()
+
+if(hsa-runtime64_FOUND AND NOT "amdgpu" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS AND NOT WIN32)
message(STATUS "Building AMDGPU plugin linked against libhsa")
target_link_libraries(omptarget.rtl.amdgpu PRIVATE hsa-runtime64::hsa-runtime64)
else()
diff --git a/offload/plugins-nextgen/cuda/CMakeLists.txt b/offload/plugins-nextgen/cuda/CMakeLists.txt
index b96e25e3dc517..b3e7b7eeaa10a 100644
--- a/offload/plugins-nextgen/cuda/CMakeLists.txt
+++ b/offload/plugins-nextgen/cuda/CMakeLists.txt
@@ -8,7 +8,13 @@ target_compile_definitions(omptarget.rtl.cuda PRIVATE OFFLOAD_MIN_CUDA_VERSION=$
find_package(CUDAToolkit QUIET ${OFFLOAD_MIN_CUDA_VERSION})
-if(CUDAToolkit_FOUND AND NOT "cuda" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS)
+if (WIN32 AND NOT "cuda" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS)
+ message(WARNING "Direct linking against libcuda is not supported on Windows; "
+ "falling back to dlopen. Add 'cuda' to "
+ "LIBOMPTARGET_DLOPEN_PLUGINS to silence this warning.")
+endif()
+
+if(CUDAToolkit_FOUND AND NOT "cuda" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS AND NOT WIN32)
message(STATUS "Building CUDA plugin linked against libcuda")
target_link_libraries(omptarget.rtl.cuda PRIVATE CUDA::cuda_driver)
else()
diff --git a/offload/plugins-nextgen/level_zero/CMakeLists.txt b/offload/plugins-nextgen/level_zero/CMakeLists.txt
index c40a375166f8b..15ca8d90ba714 100644
--- a/offload/plugins-nextgen/level_zero/CMakeLists.txt
+++ b/offload/plugins-nextgen/level_zero/CMakeLists.txt
@@ -34,7 +34,13 @@ target_include_directories(omptarget.rtl.level_zero PRIVATE
${LIBOMPTARGET_LLVM_INCLUDE_DIRS}
)
-if (LIBOMPTARGET_DEP_LEVEL_ZERO_LIBRARY AND NOT "level_zero" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS)
+if (WIN32 AND NOT "level_zero" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS)
+ message(WARNING "Direct linking against level_zero library is not supported on "
+ "Windows; falling back to dlopen. Add 'level_zero' to "
+ "LIBOMPTARGET_DLOPEN_PLUGINS to silence this warning.")
+endif()
+
+if (LIBOMPTARGET_DEP_LEVEL_ZERO_LIBRARY AND NOT "level_zero" IN_LIST LIBOMPTARGET_DLOPEN_PLUGINS AND NOT WIN32)
message(STATUS "Building Level Zero NG plugin linked against level_zero library")
target_include_directories(omptarget.rtl.level_zero PRIVATE
${LIBOMPTARGET_DEP_LEVEL_ZERO_INCLUDE_DIR}
>From b24c764378b6aaff536dfd58f29c6061312bc1b7 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Mon, 27 Jul 2026 16:29:52 +0200
Subject: [PATCH 08/20] [offload] fix missing endline
---
offload/plugins-nextgen/common/include/APIHelpers.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 182f11fc4bc3a..6e52e34793ef4 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -42,4 +42,4 @@ template <auto Fn> bool canCall() {
} // namespace api_helper
-#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_APIHELPERS_H
\ No newline at end of file
+#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_APIHELPERS_H
>From 7e213bddb77282a2b1e0d97bfb203cc6d5305968 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 29 Jul 2026 15:24:43 +0200
Subject: [PATCH 09/20] [offload] Rework `API_HELPER_OPTIONAL` so it no longer
uses UB templates
---
.../common/include/APIHelpers.h | 22 +++++++++++++------
.../plugins-nextgen/common/include/DLWrap.h | 8 +------
2 files changed, 16 insertions(+), 14 deletions(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 6e52e34793ef4..ef59442f1604c 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -17,18 +17,26 @@
#include "DLWrap.h"
// Macro to mark external symbol as weak, so linker will be okay
-// if the symbol is missing. For direct linking (dlwrap::IsDlOpened<&name> ==
-// false), we need to check if linker could find the symbol. For symbols loaded
-// using dlsym we use dlwrap::loaded<name>().
+// if the symbol is missing. For direct linking only available on Linux, we need
+// to check if linker could find the symbol. For symbols loaded using dlsym we
+// call name##_loaded function. The name##_loaded function will be nullptr if
+// external library was linked directly.
#define API_HELPER_OPTIONAL(return_type, name, ...) \
+ namespace dlwrap { \
+ bool name##_loaded() __attribute__((weak)); \
+ } \
extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
template <> inline bool api_helper::canCall<name>() { \
if (name == nullptr) \
- /* Not loaded weak symbol */ \
+ /* Not loaded weak symbol, only possible on Linux */ \
return false; \
- /* Symbols from dlwrap are never nullptr, but `loaded` might return false \
- */ \
- return dlwrap::loaded<name>(); \
+ /* If symbol wasn't dlwrapped, i.e name##_loaded == nullptr and is not \
+ * nullptr, it means the symbol was linked directly, so we can call it */ \
+ if (dlwrap::name##_loaded == nullptr) \
+ return true; \
+ /* Symbol is not nullptr and it was dlwrapped, all symbols on Windows go \
+ * here*/ \
+ return dlwrap::name##_loaded(); \
}
namespace api_helper {
diff --git a/offload/plugins-nextgen/common/include/DLWrap.h b/offload/plugins-nextgen/common/include/DLWrap.h
index 16f5a86911da3..0a33e8ef97db7 100644
--- a/offload/plugins-nextgen/common/include/DLWrap.h
+++ b/offload/plugins-nextgen/common/include/DLWrap.h
@@ -137,10 +137,6 @@ template <size_t Requested, size_t Required> constexpr void verboseAssert() {
static_assert(Requested == Required, "Arity Error");
}
-// Template to check if a symbol was loaded successfully.
-// Returns true for symbols that were not wrapped by dlwrap.
-template <auto Fn> bool loaded() { return true; }
-
} // namespace dlwrap
#define DLWRAP_INSTANTIATE(SYM_DEF, SYM_USE, ARITY) \
@@ -173,9 +169,7 @@ template <auto Fn> bool loaded() { return true; }
return reinterpret_cast<T::FunctionType>(P); \
} \
}; \
- template <> bool loaded<&(::SYMBOL)>() { \
- return SYMBOL##_Trait::get() != nullptr; \
- } \
+ bool SYMBOL##_loaded() { return SYMBOL##_Trait::get() != nullptr; } \
}
#define DLWRAP_IMPL(SYMBOL, ARITY) \
>From a8ee334281988f08a750d5ae8f3a47d6e12ffef1 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 29 Jul 2026 15:31:38 +0200
Subject: [PATCH 10/20] [offload] Make Windows path cleaner
---
.../plugins-nextgen/common/include/APIHelpers.h | 17 ++++++++++++++---
1 file changed, 14 insertions(+), 3 deletions(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index ef59442f1604c..038f94a8caa73 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -21,6 +21,7 @@
// to check if linker could find the symbol. For symbols loaded using dlsym we
// call name##_loaded function. The name##_loaded function will be nullptr if
// external library was linked directly.
+#ifndef WIN32
#define API_HELPER_OPTIONAL(return_type, name, ...) \
namespace dlwrap { \
bool name##_loaded() __attribute__((weak)); \
@@ -28,16 +29,26 @@
extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
template <> inline bool api_helper::canCall<name>() { \
if (name == nullptr) \
- /* Not loaded weak symbol, only possible on Linux */ \
+ /* Not loaded weak symbol */ \
return false; \
/* If symbol wasn't dlwrapped, i.e name##_loaded == nullptr and is not \
* nullptr, it means the symbol was linked directly, so we can call it */ \
if (dlwrap::name##_loaded == nullptr) \
return true; \
- /* Symbol is not nullptr and it was dlwrapped, all symbols on Windows go \
- * here*/ \
+ /* Symbol is not nullptr and it was dlwrapped */ \
return dlwrap::name##_loaded(); \
}
+#else
+#define API_HELPER_OPTIONAL(return_type, name, ...) \
+ namespace dlwrap { \
+ bool name##_loaded() __attribute__((weak)); \
+ } \
+ extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
+ template <> inline bool api_helper::canCall<name>() { \
+ /* Windows has no direct linking, so all symbols go through dlwrap */ \
+ return dlwrap::name##_loaded(); \
+ }
+#endif
namespace api_helper {
>From 9cb70be1a9dfa603971d27facad178edf996db8c Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 29 Jul 2026 15:46:18 +0200
Subject: [PATCH 11/20] Revert "[offload] Make Windows path cleaner"
This reverts commit 3217e64d0567da2f8a03d426bdf30e14bd753156.
---
.../plugins-nextgen/common/include/APIHelpers.h | 17 +++--------------
1 file changed, 3 insertions(+), 14 deletions(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 038f94a8caa73..ef59442f1604c 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -21,7 +21,6 @@
// to check if linker could find the symbol. For symbols loaded using dlsym we
// call name##_loaded function. The name##_loaded function will be nullptr if
// external library was linked directly.
-#ifndef WIN32
#define API_HELPER_OPTIONAL(return_type, name, ...) \
namespace dlwrap { \
bool name##_loaded() __attribute__((weak)); \
@@ -29,26 +28,16 @@
extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
template <> inline bool api_helper::canCall<name>() { \
if (name == nullptr) \
- /* Not loaded weak symbol */ \
+ /* Not loaded weak symbol, only possible on Linux */ \
return false; \
/* If symbol wasn't dlwrapped, i.e name##_loaded == nullptr and is not \
* nullptr, it means the symbol was linked directly, so we can call it */ \
if (dlwrap::name##_loaded == nullptr) \
return true; \
- /* Symbol is not nullptr and it was dlwrapped */ \
+ /* Symbol is not nullptr and it was dlwrapped, all symbols on Windows go \
+ * here*/ \
return dlwrap::name##_loaded(); \
}
-#else
-#define API_HELPER_OPTIONAL(return_type, name, ...) \
- namespace dlwrap { \
- bool name##_loaded() __attribute__((weak)); \
- } \
- extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
- template <> inline bool api_helper::canCall<name>() { \
- /* Windows has no direct linking, so all symbols go through dlwrap */ \
- return dlwrap::name##_loaded(); \
- }
-#endif
namespace api_helper {
>From 60c466570cfd2ff4dbc4b1b8c5a16cbd4d3b5ee1 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 29 Jul 2026 16:06:35 +0200
Subject: [PATCH 12/20] [offload] Remove fallback for
`zeCommandListAppendLaunchKernelWithArguments`
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 91 +------------------
1 file changed, 1 insertion(+), 90 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 22ba1106f15c8..19ecfe1c70ffa 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -110,99 +110,10 @@ 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 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) {
-
- static std::once_flag zexKernelGetArgumentSize_once;
- ze_result_t Res;
-
- // 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";
- }
- }
- }
- });
- 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 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;
-} ZeFallbacksTbl[] = {
- {"zeCommandListAppendLaunchKernelWithArguments",
- reinterpret_cast<void *>(
- &zeCommandListAppendLaunchKernelWithArgumentsFallback)}};
+} ZeFallbacksTbl[] = {};
constexpr size_t ZeFallbacksTblSz =
sizeof(ZeFallbacksTbl) / sizeof(ZeFallbacksTbl[0]);
>From 3f605714ddf45013c4944101a282a9a27f5a8c6b Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 29 Jul 2026 16:09:40 +0200
Subject: [PATCH 13/20] [offload] Allow missing symbols in dynamic level zero
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 18 ++++++++++--------
1 file changed, 10 insertions(+), 8 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 19ecfe1c70ffa..898869a8b42a8 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -170,22 +170,24 @@ static bool loadLevelZero() {
void *P = DynlibHandle->getAddressOfSymbol(Sym);
void *Fallback = nullptr;
- if (P == nullptr) {
+ if (P)
+ ODBG(OLDT_Init) << "Implementing " << Sym << " with dlsym(" << Sym
+ << ") -> " << P;
+ else {
Fallback = findZeFallback(Sym);
- if (!Fallback) {
+ if (Fallback) {
+ ODBG(OLDT_Init) << "Symbol '" << Sym << "' not found in '" << L0Library
+ << "'. Using fallback implementation -> " << Fallback;
+ P = Fallback;
+ } else {
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;
}
- if (P)
- ODBG(OLDT_Init) << "Implementing " << Sym << " with dlsym(" << Sym
- << ") -> " << P;
- *dlwrap::pointer(I) = P ? P : Fallback;
+ *dlwrap::pointer(I) = P;
}
return true;
>From 878926c6e2d06374f6d33707efb67ae84b7cc01b Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 29 Jul 2026 16:10:29 +0200
Subject: [PATCH 14/20] [offload] Add fallback without extensions in level zero
---
.../level_zero/include/L0CmdListManager.h | 6 ---
.../level_zero/src/L0Queue.cpp | 38 +++++++++++++++++--
2 files changed, 35 insertions(+), 9 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
index 1f86006b30a0b..0ce59bd8715b2 100644
--- a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
+++ b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
@@ -129,12 +129,6 @@ 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>())
- return Plugin::error(
- ErrorCode::UNSUPPORTED,
- "zeCommandListAppendLaunchKernelWithArguments is not "
- "available on this driver");
-
ze_command_list_append_launch_kernel_param_cooperative_desc_t CoopDesc = {
ZE_STRUCTURE_TYPE_COMMAND_LIST_APPEND_PARAM_COOPERATIVE_DESC, nullptr,
static_cast<ze_bool_t>(IsCooperative)};
diff --git a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
index a27024ad66d19..9d627c2646c54 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "L0Queue.h"
+#include "APIHelpers.h"
#include "L0Device.h"
#include "L0Kernel.h"
#include "L0Plugin.h"
@@ -52,9 +53,40 @@ 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);
+ auto CanUseArgPtr =
+ api_helper::canCall<zeCommandListAppendLaunchKernelWithArguments>();
+ if (KEnv.IsPtrArg && CanUseArgPtr)
+ return CmdList->appendLaunchKernelWithArgs(
+ Kernel, &KEnv.GroupCounts, &KEnv.GroupSizes, KEnv.ArgPtrs, SignalEvent,
+ NumWaitEvents, WaitEvents, KEnv.IsCooperative);
+
+ // Arguments were provided but we have an old level zero version
+ if (KEnv.IsPtrArg) {
+ 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 3ef84ea143117e7751f4f01f3d4edacf275efe2a Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Wed, 29 Jul 2026 16:57:50 +0200
Subject: [PATCH 15/20] [offload] fix APIHelpers to override existing
declarations with weak
---
offload/plugins-nextgen/common/include/APIHelpers.h | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index ef59442f1604c..7660ce332dff1 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -16,12 +16,18 @@
#include "DLWrap.h"
+#define API_HELPER_STRINGIFY_INNER(x) #x
+#define API_HELPER_STRINGIFY(x) API_HELPER_STRINGIFY_INNER(x)
+
// Macro to mark external symbol as weak, so linker will be okay
// if the symbol is missing. For direct linking only available on Linux, we need
// to check if linker could find the symbol. For symbols loaded using dlsym we
// call name##_loaded function. The name##_loaded function will be nullptr if
// external library was linked directly.
+// _Pragma("weak name") retroactively downgrades any prior strong declaration
+// (e.g. from a vendor header included before this macro) to weak.
#define API_HELPER_OPTIONAL(return_type, name, ...) \
+ _Pragma(API_HELPER_STRINGIFY(weak name)) \
namespace dlwrap { \
bool name##_loaded() __attribute__((weak)); \
} \
>From f5ec1c130c4a0d6297cc174600eba43edc2df7ba Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Thu, 30 Jul 2026 11:03:05 +0200
Subject: [PATCH 16/20] Revert "[offload] Add fallback without extensions in
level zero"
This reverts commit d3e6028b938a9e7751de7fe4913c413d7e84eb5c.
---
.../level_zero/include/L0CmdListManager.h | 6 ++++
.../level_zero/src/L0Queue.cpp | 29 +------------------
2 files changed, 7 insertions(+), 28 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
index 0ce59bd8715b2..1f86006b30a0b 100644
--- a/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
+++ b/offload/plugins-nextgen/level_zero/include/L0CmdListManager.h
@@ -129,6 +129,12 @@ 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>())
+ return Plugin::error(
+ ErrorCode::UNSUPPORTED,
+ "zeCommandListAppendLaunchKernelWithArguments is not "
+ "available on this driver");
+
ze_command_list_append_launch_kernel_param_cooperative_desc_t CoopDesc = {
ZE_STRUCTURE_TYPE_COMMAND_LIST_APPEND_PARAM_COOPERATIVE_DESC, nullptr,
static_cast<ze_bool_t>(IsCooperative)};
diff --git a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
index 9d627c2646c54..5c0ec4ad0d6c0 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
@@ -11,7 +11,6 @@
//===----------------------------------------------------------------------===//
#include "L0Queue.h"
-#include "APIHelpers.h"
#include "L0Device.h"
#include "L0Kernel.h"
#include "L0Plugin.h"
@@ -53,37 +52,11 @@ 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(); });
- auto CanUseArgPtr =
- api_helper::canCall<zeCommandListAppendLaunchKernelWithArguments>();
- if (KEnv.IsPtrArg && CanUseArgPtr)
+ if (KEnv.IsPtrArg)
return CmdList->appendLaunchKernelWithArgs(
Kernel, &KEnv.GroupCounts, &KEnv.GroupSizes, KEnv.ArgPtrs, SignalEvent,
NumWaitEvents, WaitEvents, KEnv.IsCooperative);
- // Arguments were provided but we have an old level zero version
- if (KEnv.IsPtrArg) {
- 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);
>From ea02b85b539b90d03e0c514fde5c75869251e6bb Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Thu, 30 Jul 2026 11:03:19 +0200
Subject: [PATCH 17/20] Revert "[offload] Remove fallback for
`zeCommandListAppendLaunchKernelWithArguments`"
This reverts commit 56836f50ac0a4c6e6e75dc415e9a536a51a0631f.
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 91 ++++++++++++++++++-
1 file changed, 90 insertions(+), 1 deletion(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 898869a8b42a8..09199490e4858 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -110,10 +110,99 @@ 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 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) {
+
+ static std::once_flag zexKernelGetArgumentSize_once;
+ ze_result_t Res;
+
+ // 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";
+ }
+ }
+ }
+ });
+ 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 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;
-} ZeFallbacksTbl[] = {};
+} ZeFallbacksTbl[] = {
+ {"zeCommandListAppendLaunchKernelWithArguments",
+ reinterpret_cast<void *>(
+ &zeCommandListAppendLaunchKernelWithArgumentsFallback)}};
constexpr size_t ZeFallbacksTblSz =
sizeof(ZeFallbacksTbl) / sizeof(ZeFallbacksTbl[0]);
>From b74d835ea8e87c976c41b3476460f8ca512e82fc Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Thu, 30 Jul 2026 11:26:42 +0200
Subject: [PATCH 18/20] [offload] Add `Available` check for fallbacks in
level_zero
---
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 39 ++++++++++++-------
1 file changed, 24 insertions(+), 15 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 09199490e4858..35aa10d822b10 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -114,14 +114,8 @@ DLWRAP_FINALIZE()
static ze_result_t (*zexKernelGetArgumentSize_ptr)(ze_kernel_handle_t, uint32_t,
uint32_t *) = 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) {
-
+static bool zeCommandListAppendLaunchKernelWithArgumentsFallbackAvailable() {
static std::once_flag zexKernelGetArgumentSize_once;
- ze_result_t Res;
// Load zexKernelGetArgumentSize extension if available.
std::call_once(zexKernelGetArgumentSize_once, []() {
@@ -143,13 +137,16 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
}
}
});
- 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;
- }
+ 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)
@@ -199,17 +196,29 @@ static ze_result_t zeCommandListAppendLaunchKernelWithArgumentsFallback(
static struct {
const char *Name;
void *FallbackFunc;
+ bool (*FallbackAvailable)();
} ZeFallbacksTbl[] = {
{"zeCommandListAppendLaunchKernelWithArguments",
reinterpret_cast<void *>(
- &zeCommandListAppendLaunchKernelWithArgumentsFallback)}};
+ &zeCommandListAppendLaunchKernelWithArgumentsFallback),
+ zeCommandListAppendLaunchKernelWithArgumentsFallbackAvailable}};
constexpr size_t ZeFallbacksTblSz =
sizeof(ZeFallbacksTbl) / sizeof(ZeFallbacksTbl[0]);
static void *findZeFallback(std::string_view Name) {
for (size_t i = 0; i < ZeFallbacksTblSz; i++) {
- if (Name == ZeFallbacksTbl[i].Name)
+ if (Name == ZeFallbacksTbl[i].Name) {
+ if (!ZeFallbacksTbl[i].FallbackAvailable()) {
+ ODBG(OLDT_Init)
+ << "Symbol '" << Name
+ << "' has fallback but it's not compatible with the platform!";
+ // In theory we could have multiple fallback entries for one
+ // symbol, continue the search
+ continue;
+ }
+
return ZeFallbacksTbl[i].FallbackFunc;
+ }
}
return nullptr;
}
>From 5f27afe5c3f3dcf42cc9d2a11d61c1ce87c11175 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Thu, 30 Jul 2026 11:56:47 +0200
Subject: [PATCH 19/20] [offload] fix format of APIHelpers.h
---
offload/plugins-nextgen/common/include/APIHelpers.h | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/offload/plugins-nextgen/common/include/APIHelpers.h b/offload/plugins-nextgen/common/include/APIHelpers.h
index 7660ce332dff1..30f2a10c92894 100644
--- a/offload/plugins-nextgen/common/include/APIHelpers.h
+++ b/offload/plugins-nextgen/common/include/APIHelpers.h
@@ -27,9 +27,8 @@
// _Pragma("weak name") retroactively downgrades any prior strong declaration
// (e.g. from a vendor header included before this macro) to weak.
#define API_HELPER_OPTIONAL(return_type, name, ...) \
- _Pragma(API_HELPER_STRINGIFY(weak name)) \
- namespace dlwrap { \
- bool name##_loaded() __attribute__((weak)); \
+ _Pragma(API_HELPER_STRINGIFY(weak name)) namespace dlwrap { \
+ bool name##_loaded() __attribute__((weak)); \
} \
extern "C" return_type name(__VA_ARGS__) __attribute__((weak)); \
template <> inline bool api_helper::canCall<name>() { \
>From e5e71f63b9fb164e42f1be093a23e95de5a9d6e9 Mon Sep 17 00:00:00 2001
From: blazej-smorawski <blazej.smorawski at intel.com>
Date: Thu, 30 Jul 2026 14:05:56 +0200
Subject: [PATCH 20/20] [offload] fix after rebase
---
offload/plugins-nextgen/level_zero/src/L0Queue.cpp | 11 +++--------
1 file changed, 3 insertions(+), 8 deletions(-)
diff --git a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
index 5c0ec4ad0d6c0..a27024ad66d19 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Queue.cpp
@@ -52,14 +52,9 @@ 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(); });
- if (KEnv.IsPtrArg)
- return CmdList->appendLaunchKernelWithArgs(
- Kernel, &KEnv.GroupCounts, &KEnv.GroupSizes, KEnv.ArgPtrs, SignalEvent,
- NumWaitEvents, WaitEvents, KEnv.IsCooperative);
-
- return CmdList->appendLaunchKernel(Kernel, &KEnv.GroupCounts, SignalEvent,
- NumWaitEvents, WaitEvents,
- KEnv.IsCooperative);
+ return CmdList->appendLaunchKernelWithArgs(
+ Kernel, &KEnv.GroupCounts, &KEnv.GroupSizes, KEnv.ArgPtrs, SignalEvent,
+ NumWaitEvents, WaitEvents, KEnv.IsCooperative);
}
Error L0QueueTy::memoryFill(void *Ptr, const void *Pattern, size_t PatternSize,
More information about the llvm-commits
mailing list