[llvm] [offload][l0] Report actual driver version through OL_DEVICE_INFO_DRIVER_VERSION (PR #218412)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 24 06:42:46 PDT 2026
Jan =?utf-8?q?Trusiłło?= <jan.trusillo at intel.com>
Message-ID:
In-Reply-To: <llvm.org/llvm/llvm-project/pull/218412 at github.com>
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: Jan Trusiłło (311Volt)
<details>
<summary>Changes</summary>
The Level Zero plugin currently reports the Level Zero API version as the driver version. Moreover, as the lookup table only includes API versions up to `1.12`, newer versions (e.g. `1.15`) end up being reported as Unknown.
Report the actual driver version instead:
- If the `zeIntelGetDriverVersionString` extension is available, use it. This string may be more detailed than what `zeDriverGetProperties` provides (e.g. hotfix suffix)
- If the above is not available, use `zeDriverGetProperties().driverVersion`, an U32 value whose octets make up major, minor & build numbers.
- The driver version is fetched once per context and stored in an `std::string` inside of the `L0ContextTy`.
As the PR uses L0 extensions, it uses the dispatch mechanism introduced in https://github.com/llvm/llvm-project/pull/218367 .
Stacked-on: https://github.com/llvm/llvm-project/pull/218367 . Review `9212312` and later.
Assisted-by: Codex (5.6 Sol)
---
Patch is 28.68 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/218412.diff
10 Files Affected:
- (modified) offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp (+3-88)
- (modified) offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h (+18-1)
- (modified) offload/plugins-nextgen/level_zero/include/L0CmdListManager.h (+26-8)
- (modified) offload/plugins-nextgen/level_zero/include/L0Compat.h (+17)
- (modified) offload/plugins-nextgen/level_zero/include/L0Context.h (+95-15)
- (modified) offload/plugins-nextgen/level_zero/src/L0Context.cpp (+98-16)
- (modified) offload/plugins-nextgen/level_zero/src/L0Device.cpp (+2-12)
- (modified) offload/plugins-nextgen/level_zero/src/L0Kernel.cpp (+2-2)
- (modified) offload/plugins-nextgen/level_zero/src/L0Plugin.cpp (+2-2)
- (modified) offload/plugins-nextgen/level_zero/src/L0Queue.cpp (+48-3)
``````````diff
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 55e4b6b856a96..70bc3aaf36512 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -57,6 +57,7 @@ DLWRAP(zeDeviceGetMemoryProperties, 3)
DLWRAP(zeDeviceGetCacheProperties, 3)
DLWRAP(zeDeviceGetGlobalTimestamps, 3)
DLWRAP(zeDriverGetApiVersion, 2)
+DLWRAP(zeDriverGetProperties, 2)
DLWRAP(zeDriverGetExtensionFunctionAddress, 3)
DLWRAP(zeDriverGetExtensionProperties, 3)
DLWRAP(zeEventCreate, 3)
@@ -94,6 +95,7 @@ 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..be8ad9db758ee 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
@@ -334,11 +334,24 @@ typedef struct _ze_uuid_t {
uint8_t id[16];
} ze_uuid_t;
+/* Driver UUID size */
+#ifndef ZE_MAX_DRIVER_UUID_SIZE
+#define ZE_MAX_DRIVER_UUID_SIZE 16
+#endif
+
/* Driver UUID */
typedef struct _ze_driver_uuid_t {
- ze_uuid_t id;
+ uint8_t id[ZE_MAX_DRIVER_UUID_SIZE];
} ze_driver_uuid_t;
+/* Driver properties */
+typedef struct _ze_driver_properties_t {
+ ze_structure_type_t stype;
+ void *pNext;
+ ze_driver_uuid_t uuid;
+ uint32_t driverVersion;
+} ze_driver_properties_t;
+
/* Device UUID */
typedef struct _ze_device_uuid_t {
ze_uuid_t id;
@@ -666,10 +679,14 @@ ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGet(uint32_t *pCount,
ze_driver_handle_t *phDrivers);
ZE_APIEXPORT ze_result_t ZE_APICALL
zeDriverGetApiVersion(ze_driver_handle_t hDriver, ze_api_version_t *version);
+ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGetProperties(
+ ze_driver_handle_t hDriver, ze_driver_properties_t *pDriverProperties);
ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGetExtensionFunctionAddress(
ze_driver_handle_t hDriver, const char *name, void **ppFunctionAddress);
ZE_APIEXPORT ze_result_t ZE_APICALL zeDriverGetExtensionProperties(
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..31c429b4bced9 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,27 @@ 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 +196,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..c1580991bc5b4 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Compat.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Compat.h
@@ -25,4 +25,21 @@ 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)
+
+API_HELPER_OPTIONAL(ze_result_t, zeIntelGetDriverVersionString,
+ ze_driver_handle_t hDriver, char *pDriverVersion,
+ size_t *pVersionSize)
+
#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0COMPAT_H
diff --git a/offload/plugins-nextgen/level_zero/include/L0Context.h b/offload/plugins-nextgen/level_zero/include/L0Context.h
index 09a19702baa23..7aa77d2f34247 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 setFallbackFunction(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>> {
@@ -56,6 +129,11 @@ class L0ContextTy {
/// API version supported by the Level Zero driver.
ze_api_version_t APIVersion = ZE_API_VERSION_CURRENT;
+ /// Version of the Level Zero driver.
+ std::string DriverVersion;
+
+ Expected<std::string> tryGetIntelDriverVersion();
+
/// Imported external pointers. Track this only for user-directed
/// imports/releases.
llvm::DenseMap<uintptr_t, size_t> ImportedPtrs;
@@ -133,6 +211,9 @@ class L0ContextTy {
/// Return driver API version.
ze_api_version_t getDriverAPIVersion() const { return APIVersion; }
+ /// Return driver version.
+ const std::string &getDriverVersion() const { return DriverVersion; }
+
/// Return the event pool of this driver.
EventPoolTy &getEventPool() { return EventPool; }
const EventPoolTy &getEventPool() const { return EventPool; }
@@ -145,21 +226,20 @@ 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;
+ ZeDispatcher<zeIntelGetDriverVersionString,
+ ZE_RESULT_ERROR_UNSUPPORTED_FEATURE>
+ IntelGetDriverVersionString;
};
} // namespace llvm::omp::target::plugin
diff --git a/offload/plugins-nextgen/level_zero/src/L0Context.cpp b/offload/plugins-nextgen/level_zero/src/L0Context.cpp
index cdcd1210cbcf1..6dfc0183ea2a2 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Context.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Context.cpp
@@ -13,14 +13,60 @@
#include "L0Context.h"
#include "L0Plugin.h"
+#include <cstring>
+
namespace llvm::omp::target::plugin {
+static ze_result_t ZE_APICALL getDriverVersionFromProperties(
+ ze_driver_handle_t zeDriver, char *DriverVersion, size_t *VersionSize) {
+ if (!VersionSize)
+ return ZE_RESULT_ERROR_INVALID_NULL_POINTER;
+
+ ze_driver_properties_t DriverProperties{};
+ DriverProperties.stype = ZE_STRUCTURE_TYPE_DRIVER_PROPERTIES;
+ ze_result_t Result = zeDriverGetProperties(zeDriver, &DriverProperties);
+ if (Result != ZE_RESULT_SUCCESS)
+ return Result;
+
+ uint32_t PackedVersion = DriverProperties.driverVersion;
+ std::string Version = std::to_string((PackedVersion & 0xFF000000) >> 24) +
+ "." +
+ std::to_string((PackedVersion & 0x00FF0000) >> 16) +
+ "." + std::to_string(PackedVersion & 0x0000FFFF);
+
+ if (!DriverVersion) {
+ *VersionSize = Version.size();
+ return ZE_RESULT_SUCCESS;
+ }
+
+ if (*VersionSize < Version.size()) {
+ *VersionSize = Version.size();
+ return ZE_RESULT_ERROR_INVALID_SIZE;
+ }
+
+ std::memcpy(DriverVersion, Version.data(), Version.size());
+ *VersionSize = Version.size();
+ return ZE_RESULT_SUCCESS;
+}
+
L0ContextTy::L0ContextTy(LevelZeroPluginTy &Plugin, ze_driver_handle_t zeDriver,
int32_t DriverId)
: Plugin(Plugin), zeDriver(zeDriver) {}
L0ContextTy::~L0ContextTy() = default;
+Expected<std::string> L0ContextTy::tryGetIntelDriverVersion() {
+ size_t VersionSize = 0;
+ CALL_ZE_RET_ERROR(IntelGetDriverVersionString, zeDriver, nullptr,
+ &VersionSize);
+ std::string Version(VersionSize, '\0');
+ CALL_ZE_RET_ERROR(IntelGetDriverVersionString, zeDriver, Version.data(),
+ &VersionSize);
+ if (!Version.empty() && Version.back() == '\0')
+ Version.pop_back();
+ return Version;
+}
+
Error L0ContextTy::init() {
auto CleanupOnError = [&]() {
if (zeContext) {
@@ -57,22 +103,58 @@ 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...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/218412
More information about the llvm-commits
mailing list