diff --git a/hsa.cpp b/hsa.cpp new file mode 100755 index 0000000000..a596217598 --- /dev/null +++ b/hsa.cpp @@ -0,0 +1,120 @@ +#include +#include "libhsakmt.h" +#include "hsa-runtime/inc/hsa.h" +#include "hsa-runtime/inc/hsa_ven_amd_loader.h" + +static std::unique_ptr lock_ = std::make_unique(); +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 gard(*lock_); \ + if (_sym == nullptr) { \ + fn_##_sym = \ + reinterpret_cast(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 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(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 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 diff --git a/openclose.cpp b/openclose.cpp index e30f17d6d4..0783f401ea 100644 --- a/openclose.cpp +++ b/openclose.cpp @@ -30,22 +30,9 @@ #include #include #include -#include #include #include #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(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; diff --git a/wddm/queue.cpp b/wddm/queue.cpp index 28fcd40ffc..5ffaa12e31 100644 --- a/wddm/queue.cpp +++ b/wddm/queue.cpp @@ -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(addr), + auto ret = hsakmt_hsa_ven_amd_loader_query_host_address(reinterpret_cast(addr), reinterpret_cast(&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;