wsl/hsakmt: drop dlsym for hsa symbols
Signed-off-by: Flora Cui <flora.cui@amd.com>
Αυτή η υποβολή περιλαμβάνεται σε:
Flora Cui
Frank Min
Εκτελέσιμο αρχείο
+120
@@ -0,0 +1,120 @@
|
||||
#include <dlfcn.h>
|
||||
#include "libhsakmt.h"
|
||||
#include "hsa-runtime/inc/hsa.h"
|
||||
#include "hsa-runtime/inc/hsa_ven_amd_loader.h"
|
||||
|
||||
static std::unique_ptr<std::mutex> lock_ = std::make_unique<std::mutex>();
|
||||
static hsa_status_t (*fn_hsa_ven_amd_loader_query_host_address)(
|
||||
const void *device_address, const void **host_address);
|
||||
|
||||
#if 0
|
||||
static hsa_signal_value_t (*fn_hsa_signal_load_relaxed)(hsa_signal_t signal);
|
||||
static hsa_signal_value_t (*fn_hsa_signal_wait_relaxed)(
|
||||
hsa_signal_t signal, hsa_signal_condition_t condition,
|
||||
hsa_signal_value_t compare_value, uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_state_hint);
|
||||
static void (*fn_hsa_signal_store_screlease)(hsa_signal_t hsa_signal,
|
||||
hsa_signal_value_t value);
|
||||
|
||||
#define _HSAKMT_LOOKUP_SYMS(_sym) \
|
||||
if (_sym == nullptr) { \
|
||||
std::lock_guard<std::mutex> gard(*lock_); \
|
||||
if (_sym == nullptr) { \
|
||||
fn_##_sym = \
|
||||
reinterpret_cast<decltype(fn_##_sym)>(dlsym(RTLD_DEFAULT, #_sym)); \
|
||||
if (!fn_##_sym) { \
|
||||
pr_err("%s not found - %s\n", #_sym, dlerror()); \
|
||||
} \
|
||||
} \
|
||||
}
|
||||
|
||||
#define _HSAKMT_EXEC_API(_sym, ...) \
|
||||
do { \
|
||||
if (fn_##_sym != nullptr) { \
|
||||
return fn_##_sym(__VA_ARGS__); \
|
||||
} \
|
||||
} while(0);
|
||||
|
||||
hsa_signal_value_t hsakmt_hsa_signal_load_relaxed(hsa_signal_t signal) {
|
||||
_HSAKMT_LOOKUP_SYMS(hsa_signal_load_relaxed);
|
||||
_HSAKMT_EXEC_API(hsa_signal_load_relaxed, signal);
|
||||
return 0;
|
||||
}
|
||||
|
||||
hsa_signal_value_t hsakmt_hsa_signal_wait_relaxed(
|
||||
hsa_signal_t signal, hsa_signal_condition_t condition,
|
||||
hsa_signal_value_t compare_value, uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_state_hint) {
|
||||
_HSAKMT_LOOKUP_SYMS(hsa_signal_wait_relaxed);
|
||||
_HSAKMT_EXEC_API(hsa_signal_wait_relaxed, signal, condition, compare_value,
|
||||
timeout_hint, wait_state_hint);
|
||||
return 0;
|
||||
}
|
||||
|
||||
void hsakmt_hsa_signal_store_screlease(hsa_signal_t hsa_signal,
|
||||
hsa_signal_value_t value){
|
||||
_HSAKMT_LOOKUP_SYMS(hsa_signal_store_screlease);
|
||||
_HSAKMT_EXEC_API(hsa_signal_store_screlease, hsa_signal, value);
|
||||
}
|
||||
|
||||
hsa_status_t hsakmt_hsa_ven_amd_loader_query_host_address(
|
||||
const void *device_address, const void **host_address) {
|
||||
if (fn_hsa_ven_amd_loader_query_host_address == nullptr) {
|
||||
std::lock_guard<std::mutex> gard(*lock_);
|
||||
if (fn_hsa_ven_amd_loader_query_host_address == nullptr) {
|
||||
hsa_status_t (*fn_hsa_system_get_extension_table)(
|
||||
uint16_t extension, uint16_t version_major, uint16_t version_minor, void *table);
|
||||
fn_hsa_system_get_extension_table =
|
||||
reinterpret_cast<decltype(fn_hsa_system_get_extension_table)>(dlsym(RTLD_DEFAULT, "hsa_system_get_extension_table"));
|
||||
if (fn_hsa_system_get_extension_table == nullptr) {
|
||||
pr_err("%s not found - %s\n", "hsa_system_get_extension_table", dlerror());
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
|
||||
hsa_ven_amd_loader_1_03_pfn_t table;
|
||||
fn_hsa_system_get_extension_table(HSA_EXTENSION_AMD_LOADER, 1, 3, &table);
|
||||
fn_hsa_ven_amd_loader_query_host_address =
|
||||
table.hsa_ven_amd_loader_query_host_address;
|
||||
}
|
||||
}
|
||||
|
||||
_HSAKMT_EXEC_API(hsa_ven_amd_loader_query_host_address, device_address, host_address);
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
|
||||
#else
|
||||
hsa_signal_value_t hsakmt_hsa_signal_load_relaxed(hsa_signal_t signal) {
|
||||
return hsa_signal_load_relaxed(signal);
|
||||
}
|
||||
|
||||
hsa_signal_value_t hsakmt_hsa_signal_wait_relaxed(
|
||||
hsa_signal_t signal, hsa_signal_condition_t condition,
|
||||
hsa_signal_value_t compare_value, uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_state_hint) {
|
||||
return hsa_signal_wait_relaxed(signal, condition, compare_value, timeout_hint,
|
||||
wait_state_hint);
|
||||
}
|
||||
|
||||
void hsakmt_hsa_signal_store_screlease(hsa_signal_t hsa_signal,
|
||||
hsa_signal_value_t value) {
|
||||
hsa_signal_store_screlease(hsa_signal, value);
|
||||
}
|
||||
|
||||
hsa_status_t hsakmt_hsa_ven_amd_loader_query_host_address(
|
||||
const void *device_address, const void **host_address) {
|
||||
if (fn_hsa_ven_amd_loader_query_host_address == nullptr) {
|
||||
std::lock_guard<std::mutex> gard(*lock_);
|
||||
if (fn_hsa_ven_amd_loader_query_host_address == nullptr) {
|
||||
hsa_ven_amd_loader_1_03_pfn_t table;
|
||||
hsa_system_get_extension_table(HSA_EXTENSION_AMD_LOADER, 1, 3, &table);
|
||||
fn_hsa_ven_amd_loader_query_host_address =
|
||||
table.hsa_ven_amd_loader_query_host_address;
|
||||
}
|
||||
}
|
||||
|
||||
if (fn_hsa_ven_amd_loader_query_host_address)
|
||||
return fn_hsa_ven_amd_loader_query_host_address(device_address, host_address);
|
||||
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
#endif
|
||||
@@ -30,22 +30,9 @@
|
||||
#include <unistd.h>
|
||||
#include <cstdio>
|
||||
#include <strings.h>
|
||||
#include <dlfcn.h>
|
||||
#include <cassert>
|
||||
#include <amdgpu.h>
|
||||
#include "libhsakmt.h"
|
||||
#include "inc/hsa/hsa.h"
|
||||
#include "inc/hsa/hsa_ven_amd_loader.h"
|
||||
|
||||
hsa_signal_value_t (*fn_hsa_signal_load_relaxed)(hsa_signal_t signal);
|
||||
hsa_signal_value_t (*fn_hsa_signal_wait_relaxed)(
|
||||
hsa_signal_t signal, hsa_signal_condition_t condition,
|
||||
hsa_signal_value_t compare_value, uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_state_hint);
|
||||
void (*fn_hsa_signal_store_screlease)(hsa_signal_t hsa_signal,
|
||||
hsa_signal_value_t value);
|
||||
hsa_status_t (*fn_hsa_ven_amd_loader_query_host_address)(
|
||||
const void *device_address, const void **host_address);
|
||||
|
||||
static const char dxg_device_name[] = "/dev/dxg";
|
||||
static pid_t parent_pid = -1;
|
||||
@@ -146,33 +133,6 @@ static HSAKMT_STATUS init_vars_from_env(void) {
|
||||
return HSAKMT_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
#define _HSAKMT_LOOKUP_SYMS(_sym) \
|
||||
do { \
|
||||
fn_##_sym = \
|
||||
reinterpret_cast<decltype(fn_##_sym)>(dlsym(RTLD_DEFAULT, #_sym)); \
|
||||
if (!fn_##_sym) { \
|
||||
pr_err("%s not found - %s\n", #_sym, dlerror()); \
|
||||
return HSAKMT_STATUS_ERROR; \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
static HSAKMT_STATUS init_symbols(void) {
|
||||
_HSAKMT_LOOKUP_SYMS(hsa_signal_load_relaxed);
|
||||
_HSAKMT_LOOKUP_SYMS(hsa_signal_wait_relaxed);
|
||||
_HSAKMT_LOOKUP_SYMS(hsa_signal_store_screlease);
|
||||
|
||||
hsa_status_t (*fn_hsa_system_get_extension_table)(
|
||||
uint16_t extension, uint16_t version_major, uint16_t version_minor,
|
||||
void *table);
|
||||
_HSAKMT_LOOKUP_SYMS(hsa_system_get_extension_table);
|
||||
hsa_ven_amd_loader_1_03_pfn_t table;
|
||||
fn_hsa_system_get_extension_table(HSA_EXTENSION_AMD_LOADER, 1, 3, &table);
|
||||
fn_hsa_ven_amd_loader_query_host_address =
|
||||
table.hsa_ven_amd_loader_query_host_address;
|
||||
|
||||
return HSAKMT_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
HSAKMT_STATUS HSAKMTAPI hsaKmtOpenKFD(void) {
|
||||
HSAKMT_STATUS result;
|
||||
int fd = -1;
|
||||
@@ -192,10 +152,6 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtOpenKFD(void) {
|
||||
if (dxg_open_count == 0) {
|
||||
static bool atfork_installed = false;
|
||||
|
||||
result = init_symbols();
|
||||
if (result != HSAKMT_STATUS_SUCCESS)
|
||||
goto open_failed;
|
||||
|
||||
result = init_vars_from_env();
|
||||
if (result != HSAKMT_STATUS_SUCCESS)
|
||||
goto open_failed;
|
||||
|
||||
+18
-18
@@ -48,6 +48,18 @@
|
||||
#include "inc/registers.h"
|
||||
#include "libhsakmt.h"
|
||||
|
||||
#include "hsa-runtime/inc/hsa.h"
|
||||
#include "hsa-runtime/inc/hsa_ven_amd_loader.h"
|
||||
extern hsa_signal_value_t hsakmt_hsa_signal_load_relaxed(hsa_signal_t signal);
|
||||
extern hsa_signal_value_t hsakmt_hsa_signal_wait_relaxed(
|
||||
hsa_signal_t signal, hsa_signal_condition_t condition,
|
||||
hsa_signal_value_t compare_value, uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_state_hint);
|
||||
extern void hsakmt_hsa_signal_store_screlease(hsa_signal_t hsa_signal,
|
||||
hsa_signal_value_t value);
|
||||
extern hsa_status_t hsakmt_hsa_ven_amd_loader_query_host_address(
|
||||
const void *device_address, const void **host_address);
|
||||
|
||||
namespace wsl {
|
||||
namespace thunk {
|
||||
|
||||
@@ -130,7 +142,6 @@ hsa_status_t WDDMQueue::SetPriority(hsa_amd_queue_priority_t priority) {
|
||||
return HwsInit();
|
||||
}
|
||||
|
||||
extern "C" void (*fn_hsa_signal_store_screlease)(hsa_signal_t hsa_signal, hsa_signal_value_t value);
|
||||
void ComputeQueue::HandleError(hsa_status_t status) {
|
||||
hsa_signal_t sig = amd_queue_rocr_->queue_inactive_signal;
|
||||
hsa_signal_value_t val = -1;
|
||||
@@ -160,7 +171,7 @@ void ComputeQueue::HandleError(hsa_status_t status) {
|
||||
}
|
||||
|
||||
if (sig.handle) {
|
||||
fn_hsa_signal_store_screlease(sig, val);
|
||||
hsakmt_hsa_signal_store_screlease(sig, val);
|
||||
}
|
||||
if (error_code_) {
|
||||
error_code_->store(val, std::memory_order_release);
|
||||
@@ -505,13 +516,10 @@ uint32_t ComputeQueue::UpdateIndexStride(uint32_t srd, bool wave32) {
|
||||
return srd;
|
||||
}
|
||||
|
||||
extern "C" hsa_status_t (*fn_hsa_ven_amd_loader_query_host_address)(
|
||||
const void *device_address,
|
||||
const void **host_address);
|
||||
uint64_t ComputeQueue::GetKernelObjAddr(uint64_t addr) const {
|
||||
//TODO: convert dev_addr to host_addr
|
||||
uint64_t host_addr = 0;
|
||||
auto ret = fn_hsa_ven_amd_loader_query_host_address(reinterpret_cast<const void *>(addr),
|
||||
auto ret = hsakmt_hsa_ven_amd_loader_query_host_address(reinterpret_cast<const void *>(addr),
|
||||
reinterpret_cast<const void **>(&host_addr));
|
||||
if (ret == HSA_STATUS_ERROR_INVALID_ARGUMENT) {
|
||||
return NULL;
|
||||
@@ -706,14 +714,6 @@ ComputeQueue::KernelDispatchAqlToPm4(char *cpu, hsa_kernel_dispatch_packet_t *pa
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
extern "C" hsa_signal_value_t (*fn_hsa_signal_load_relaxed)(
|
||||
hsa_signal_t signal);
|
||||
extern "C" hsa_signal_value_t (*fn_hsa_signal_wait_relaxed)(
|
||||
hsa_signal_t signal,
|
||||
hsa_signal_condition_t condition,
|
||||
hsa_signal_value_t compare_value,
|
||||
uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_state_hint);
|
||||
hsa_status_t
|
||||
ComputeQueue::BarrierGenericAqlToPm4(char *cpu, hsa_barrier_and_packet_t *packet, bool is_or) {
|
||||
debug_print("queue %p %s head=%x dep %" PRIx64 " %" PRIx64 " %" PRIx64
|
||||
@@ -735,7 +735,7 @@ ComputeQueue::BarrierGenericAqlToPm4(char *cpu, hsa_barrier_and_packet_t *packet
|
||||
|
||||
while (n) {
|
||||
for (int i = 0; i < n; i++) {
|
||||
if (!fn_hsa_signal_load_relaxed(sig[i])) {
|
||||
if (!hsakmt_hsa_signal_load_relaxed(sig[i])) {
|
||||
unsignaled = false;
|
||||
break;
|
||||
}
|
||||
@@ -751,7 +751,7 @@ ComputeQueue::BarrierGenericAqlToPm4(char *cpu, hsa_barrier_and_packet_t *packet
|
||||
continue;
|
||||
|
||||
hsa_signal_value_t value =
|
||||
fn_hsa_signal_wait_relaxed(packet->dep_signal[i], HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
|
||||
hsakmt_hsa_signal_wait_relaxed(packet->dep_signal[i], HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
|
||||
assert(value == 0);
|
||||
}
|
||||
}
|
||||
@@ -851,7 +851,7 @@ hsa_status_t ComputeQueue::VendorSpecificAqlToPm4(char *cpu, amd_aql_pm4_ib *pac
|
||||
ib_size = i;
|
||||
} else {
|
||||
if (packet->completion_signal.handle != 0) {
|
||||
fn_hsa_signal_store_screlease(packet->completion_signal, 0);
|
||||
hsakmt_hsa_signal_store_screlease(packet->completion_signal, 0);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -986,7 +986,7 @@ void SDMAQueue::SdmaThread(SDMAQueue *queue) {
|
||||
debug_print("SDMA: poll signal %#lx addr %#lx val %d\n", signal_handle, poll_addr, poll_val);
|
||||
hsa_signal_t hsa_signal = {signal_handle};
|
||||
hsa_signal_value_t value =
|
||||
fn_hsa_signal_wait_relaxed(hsa_signal, HSA_SIGNAL_CONDITION_EQ, poll_val, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
|
||||
hsakmt_hsa_signal_wait_relaxed(hsa_signal, HSA_SIGNAL_CONDITION_EQ, poll_val, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
|
||||
assert(value == poll_val);
|
||||
|
||||
poll_pkt += 2;
|
||||
|
||||
Αναφορά σε νέο ζήτημα
Block a user