Export Amd Extension APIs including support for Version Control
Change-Id: I8c03cbd4049e8115ae00d51f193b9c31ac941f21
[ROCm/ROCR-Runtime commit: 95dc97da7b]
Dieser Commit ist enthalten in:
Ramesh Errabolu
+324
-99
@@ -41,60 +41,69 @@
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
#include "hsa_api_trace.h"
|
||||
#include "core/inc/hsa_api_trace_int.h"
|
||||
|
||||
static const ApiTable* HsaApiTable;
|
||||
static const HsaApiTable* hsaApiTable;
|
||||
static const CoreApiTable* coreApiTable;
|
||||
static const AmdExtTable* amdExtTable;
|
||||
|
||||
void hsa_table_interface_init(const ApiTable* Table) { HsaApiTable = Table; }
|
||||
void hsa_table_interface_init(const HsaApiTable* apiTable) {
|
||||
hsaApiTable = apiTable;
|
||||
coreApiTable = apiTable->core_;
|
||||
amdExtTable = apiTable->amd_ext_;
|
||||
}
|
||||
|
||||
const ApiTable* hsa_table_interface_get_table() { return HsaApiTable; }
|
||||
const HsaApiTable* hsa_table_interface_get_table() {
|
||||
return hsaApiTable;
|
||||
}
|
||||
|
||||
// Pass through stub functions
|
||||
hsa_status_t HSA_API hsa_init() { return HsaApiTable->hsa_init_fn(); }
|
||||
hsa_status_t HSA_API hsa_init() { return coreApiTable->hsa_init_fn(); }
|
||||
|
||||
hsa_status_t HSA_API hsa_shut_down() { return HsaApiTable->hsa_shut_down_fn(); }
|
||||
hsa_status_t HSA_API hsa_shut_down() { return coreApiTable->hsa_shut_down_fn(); }
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_system_get_info(hsa_system_info_t attribute, void* value) {
|
||||
return HsaApiTable->hsa_system_get_info_fn(attribute, value);
|
||||
return coreApiTable->hsa_system_get_info_fn(attribute, value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_system_extension_supported(uint16_t extension, uint16_t version_major,
|
||||
uint16_t version_minor, bool* result) {
|
||||
return HsaApiTable->hsa_system_extension_supported_fn(
|
||||
return coreApiTable->hsa_system_extension_supported_fn(
|
||||
extension, version_major, version_minor, result);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_system_get_extension_table(uint16_t extension, uint16_t version_major,
|
||||
uint16_t version_minor, void* table) {
|
||||
return HsaApiTable->hsa_system_get_extension_table_fn(
|
||||
return coreApiTable->hsa_system_get_extension_table_fn(
|
||||
extension, version_major, version_minor, table);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, void* data),
|
||||
void* data) {
|
||||
return HsaApiTable->hsa_iterate_agents_fn(callback, data);
|
||||
return coreApiTable->hsa_iterate_agents_fn(callback, data);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_agent_get_info(hsa_agent_t agent,
|
||||
hsa_agent_info_t attribute,
|
||||
void* value) {
|
||||
return HsaApiTable->hsa_agent_get_info_fn(agent, attribute, value);
|
||||
return coreApiTable->hsa_agent_get_info_fn(agent, attribute, value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_agent_get_exception_policies(hsa_agent_t agent,
|
||||
hsa_profile_t profile,
|
||||
uint16_t* mask) {
|
||||
return HsaApiTable->hsa_agent_get_exception_policies_fn(agent, profile, mask);
|
||||
return coreApiTable->hsa_agent_get_exception_policies_fn(agent, profile, mask);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_agent_extension_supported(uint16_t extension, hsa_agent_t agent,
|
||||
uint16_t version_major,
|
||||
uint16_t version_minor, bool* result) {
|
||||
return HsaApiTable->hsa_agent_extension_supported_fn(
|
||||
return coreApiTable->hsa_agent_extension_supported_fn(
|
||||
extension, agent, version_major, version_minor, result);
|
||||
}
|
||||
|
||||
@@ -104,7 +113,7 @@ hsa_status_t HSA_API
|
||||
void* data),
|
||||
void* data, uint32_t private_segment_size,
|
||||
uint32_t group_segment_size, hsa_queue_t** queue) {
|
||||
return HsaApiTable->hsa_queue_create_fn(agent, size, type, callback, data,
|
||||
return coreApiTable->hsa_queue_create_fn(agent, size, type, callback, data,
|
||||
private_segment_size,
|
||||
group_segment_size, queue);
|
||||
}
|
||||
@@ -113,167 +122,167 @@ hsa_status_t HSA_API
|
||||
hsa_soft_queue_create(hsa_region_t region, uint32_t size,
|
||||
hsa_queue_type_t type, uint32_t features,
|
||||
hsa_signal_t completion_signal, hsa_queue_t** queue) {
|
||||
return HsaApiTable->hsa_soft_queue_create_fn(region, size, type, features,
|
||||
return coreApiTable->hsa_soft_queue_create_fn(region, size, type, features,
|
||||
completion_signal, queue);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_queue_destroy(hsa_queue_t* queue) {
|
||||
return HsaApiTable->hsa_queue_destroy_fn(queue);
|
||||
return coreApiTable->hsa_queue_destroy_fn(queue);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_queue_inactivate(hsa_queue_t* queue) {
|
||||
return HsaApiTable->hsa_queue_inactivate_fn(queue);
|
||||
return coreApiTable->hsa_queue_inactivate_fn(queue);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_load_read_index_acquire(const hsa_queue_t* queue) {
|
||||
return HsaApiTable->hsa_queue_load_read_index_acquire_fn(queue);
|
||||
return coreApiTable->hsa_queue_load_read_index_acquire_fn(queue);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_load_read_index_relaxed(const hsa_queue_t* queue) {
|
||||
return HsaApiTable->hsa_queue_load_read_index_relaxed_fn(queue);
|
||||
return coreApiTable->hsa_queue_load_read_index_relaxed_fn(queue);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_load_write_index_acquire(const hsa_queue_t* queue) {
|
||||
return HsaApiTable->hsa_queue_load_write_index_acquire_fn(queue);
|
||||
return coreApiTable->hsa_queue_load_write_index_acquire_fn(queue);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_load_write_index_relaxed(const hsa_queue_t* queue) {
|
||||
return HsaApiTable->hsa_queue_load_write_index_relaxed_fn(queue);
|
||||
return coreApiTable->hsa_queue_load_write_index_relaxed_fn(queue);
|
||||
}
|
||||
|
||||
void HSA_API hsa_queue_store_write_index_relaxed(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_store_write_index_relaxed_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_store_write_index_relaxed_fn(queue, value);
|
||||
}
|
||||
|
||||
void HSA_API hsa_queue_store_write_index_release(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_store_write_index_release_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_store_write_index_release_fn(queue, value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_cas_write_index_acq_rel(const hsa_queue_t* queue,
|
||||
uint64_t expected,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_cas_write_index_acq_rel_fn(queue, expected,
|
||||
return coreApiTable->hsa_queue_cas_write_index_acq_rel_fn(queue, expected,
|
||||
value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_cas_write_index_acquire(const hsa_queue_t* queue,
|
||||
uint64_t expected,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_cas_write_index_acquire_fn(queue, expected,
|
||||
return coreApiTable->hsa_queue_cas_write_index_acquire_fn(queue, expected,
|
||||
value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_cas_write_index_relaxed(const hsa_queue_t* queue,
|
||||
uint64_t expected,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_cas_write_index_relaxed_fn(queue, expected,
|
||||
return coreApiTable->hsa_queue_cas_write_index_relaxed_fn(queue, expected,
|
||||
value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_cas_write_index_release(const hsa_queue_t* queue,
|
||||
uint64_t expected,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_cas_write_index_release_fn(queue, expected,
|
||||
return coreApiTable->hsa_queue_cas_write_index_release_fn(queue, expected,
|
||||
value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_add_write_index_acq_rel(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_add_write_index_acq_rel_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_add_write_index_acq_rel_fn(queue, value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_add_write_index_acquire(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_add_write_index_acquire_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_add_write_index_acquire_fn(queue, value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_add_write_index_relaxed(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_add_write_index_relaxed_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_add_write_index_relaxed_fn(queue, value);
|
||||
}
|
||||
|
||||
uint64_t HSA_API hsa_queue_add_write_index_release(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_add_write_index_release_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_add_write_index_release_fn(queue, value);
|
||||
}
|
||||
|
||||
void HSA_API hsa_queue_store_read_index_relaxed(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_store_read_index_relaxed_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_store_read_index_relaxed_fn(queue, value);
|
||||
}
|
||||
|
||||
void HSA_API hsa_queue_store_read_index_release(const hsa_queue_t* queue,
|
||||
uint64_t value) {
|
||||
return HsaApiTable->hsa_queue_store_read_index_release_fn(queue, value);
|
||||
return coreApiTable->hsa_queue_store_read_index_release_fn(queue, value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_agent_iterate_regions(
|
||||
hsa_agent_t agent,
|
||||
hsa_status_t (*callback)(hsa_region_t region, void* data), void* data) {
|
||||
return HsaApiTable->hsa_agent_iterate_regions_fn(agent, callback, data);
|
||||
return coreApiTable->hsa_agent_iterate_regions_fn(agent, callback, data);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_region_get_info(hsa_region_t region,
|
||||
hsa_region_info_t attribute,
|
||||
void* value) {
|
||||
return HsaApiTable->hsa_region_get_info_fn(region, attribute, value);
|
||||
return coreApiTable->hsa_region_get_info_fn(region, attribute, value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_memory_register(void* address, size_t size) {
|
||||
return HsaApiTable->hsa_memory_register_fn(address, size);
|
||||
return coreApiTable->hsa_memory_register_fn(address, size);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_memory_deregister(void* address, size_t size) {
|
||||
return HsaApiTable->hsa_memory_deregister_fn(address, size);
|
||||
return coreApiTable->hsa_memory_deregister_fn(address, size);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_memory_allocate(hsa_region_t region, size_t size, void** ptr) {
|
||||
return HsaApiTable->hsa_memory_allocate_fn(region, size, ptr);
|
||||
return coreApiTable->hsa_memory_allocate_fn(region, size, ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_memory_free(void* ptr) {
|
||||
return HsaApiTable->hsa_memory_free_fn(ptr);
|
||||
return coreApiTable->hsa_memory_free_fn(ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_memory_copy(void* dst, const void* src, size_t size) {
|
||||
return HsaApiTable->hsa_memory_copy_fn(dst, src, size);
|
||||
return coreApiTable->hsa_memory_copy_fn(dst, src, size);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_memory_assign_agent(void* ptr, hsa_agent_t agent,
|
||||
hsa_access_permission_t access) {
|
||||
return HsaApiTable->hsa_memory_assign_agent_fn(ptr, agent, access);
|
||||
return coreApiTable->hsa_memory_assign_agent_fn(ptr, agent, access);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers,
|
||||
const hsa_agent_t* consumers, hsa_signal_t* signal) {
|
||||
return HsaApiTable->hsa_signal_create_fn(initial_value, num_consumers,
|
||||
return coreApiTable->hsa_signal_create_fn(initial_value, num_consumers,
|
||||
consumers, signal);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_signal_destroy(hsa_signal_t signal) {
|
||||
return HsaApiTable->hsa_signal_destroy_fn(signal);
|
||||
return coreApiTable->hsa_signal_destroy_fn(signal);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API hsa_signal_load_relaxed(hsa_signal_t signal) {
|
||||
return HsaApiTable->hsa_signal_load_relaxed_fn(signal);
|
||||
return coreApiTable->hsa_signal_load_relaxed_fn(signal);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API hsa_signal_load_acquire(hsa_signal_t signal) {
|
||||
return HsaApiTable->hsa_signal_load_acquire_fn(signal);
|
||||
return coreApiTable->hsa_signal_load_acquire_fn(signal);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_store_relaxed_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_store_relaxed_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_store_release(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_store_release_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_store_release_fn(signal, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API
|
||||
@@ -282,7 +291,7 @@ hsa_signal_value_t HSA_API
|
||||
hsa_signal_value_t compare_value,
|
||||
uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_expectancy_hint) {
|
||||
return HsaApiTable->hsa_signal_wait_relaxed_fn(
|
||||
return coreApiTable->hsa_signal_wait_relaxed_fn(
|
||||
signal, condition, compare_value, timeout_hint, wait_expectancy_hint);
|
||||
}
|
||||
|
||||
@@ -292,166 +301,166 @@ hsa_signal_value_t HSA_API
|
||||
hsa_signal_value_t compare_value,
|
||||
uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_expectancy_hint) {
|
||||
return HsaApiTable->hsa_signal_wait_acquire_fn(
|
||||
return coreApiTable->hsa_signal_wait_acquire_fn(
|
||||
signal, condition, compare_value, timeout_hint, wait_expectancy_hint);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_and_relaxed(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_and_relaxed_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_and_relaxed_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_and_acquire(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_and_acquire_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_and_acquire_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_and_release(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_and_release_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_and_release_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_and_acq_rel(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_and_acq_rel_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_and_acq_rel_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_or_relaxed(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_or_relaxed_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_or_relaxed_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_or_acquire(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_or_acquire_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_or_acquire_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_or_release(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_or_release_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_or_release_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_or_acq_rel(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_or_acq_rel_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_or_acq_rel_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_xor_relaxed(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_xor_relaxed_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_xor_relaxed_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_xor_acquire(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_xor_acquire_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_xor_acquire_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_xor_release(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_xor_release_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_xor_release_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_xor_acq_rel(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_xor_acq_rel_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_xor_acq_rel_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_add_relaxed(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_add_relaxed_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_add_relaxed_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_add_acquire(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_add_acquire_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_add_acquire_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_add_release(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_add_release_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_add_release_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_add_acq_rel(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_add_acq_rel_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_add_acq_rel_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_subtract_relaxed(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_subtract_relaxed_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_subtract_relaxed_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_subtract_acquire(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_subtract_acquire_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_subtract_acquire_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_subtract_release(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_subtract_release_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_subtract_release_fn(signal, value);
|
||||
}
|
||||
|
||||
void HSA_API
|
||||
hsa_signal_subtract_acq_rel(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_subtract_acq_rel_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_subtract_acq_rel_fn(signal, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API
|
||||
hsa_signal_exchange_relaxed(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_exchange_relaxed_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_exchange_relaxed_fn(signal, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API
|
||||
hsa_signal_exchange_acquire(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_exchange_acquire_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_exchange_acquire_fn(signal, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API
|
||||
hsa_signal_exchange_release(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_exchange_release_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_exchange_release_fn(signal, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API
|
||||
hsa_signal_exchange_acq_rel(hsa_signal_t signal, hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_exchange_acq_rel_fn(signal, value);
|
||||
return coreApiTable->hsa_signal_exchange_acq_rel_fn(signal, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API hsa_signal_cas_relaxed(hsa_signal_t signal,
|
||||
hsa_signal_value_t expected,
|
||||
hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_cas_relaxed_fn(signal, expected, value);
|
||||
return coreApiTable->hsa_signal_cas_relaxed_fn(signal, expected, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API hsa_signal_cas_acquire(hsa_signal_t signal,
|
||||
hsa_signal_value_t expected,
|
||||
hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_cas_acquire_fn(signal, expected, value);
|
||||
return coreApiTable->hsa_signal_cas_acquire_fn(signal, expected, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API hsa_signal_cas_release(hsa_signal_t signal,
|
||||
hsa_signal_value_t expected,
|
||||
hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_cas_release_fn(signal, expected, value);
|
||||
return coreApiTable->hsa_signal_cas_release_fn(signal, expected, value);
|
||||
}
|
||||
|
||||
hsa_signal_value_t HSA_API hsa_signal_cas_acq_rel(hsa_signal_t signal,
|
||||
hsa_signal_value_t expected,
|
||||
hsa_signal_value_t value) {
|
||||
return HsaApiTable->hsa_signal_cas_acq_rel_fn(signal, expected, value);
|
||||
return coreApiTable->hsa_signal_cas_acq_rel_fn(signal, expected, value);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_isa_from_name(const char* name, hsa_isa_t* isa) {
|
||||
return HsaApiTable->hsa_isa_from_name_fn(name, isa);
|
||||
return coreApiTable->hsa_isa_from_name_fn(name, isa);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_isa_get_info(hsa_isa_t isa, hsa_isa_info_t attribute,
|
||||
uint32_t index, void* value) {
|
||||
return HsaApiTable->hsa_isa_get_info_fn(isa, attribute, index, value);
|
||||
return coreApiTable->hsa_isa_get_info_fn(isa, attribute, index, value);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_isa_compatible(hsa_isa_t code_object_isa, hsa_isa_t agent_isa,
|
||||
bool* result) {
|
||||
return HsaApiTable->hsa_isa_compatible_fn(code_object_isa, agent_isa, result);
|
||||
return coreApiTable->hsa_isa_compatible_fn(code_object_isa, agent_isa, result);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_code_object_serialize(
|
||||
@@ -460,7 +469,7 @@ hsa_status_t HSA_API hsa_code_object_serialize(
|
||||
void** address),
|
||||
hsa_callback_data_t callback_data, const char* options,
|
||||
void** serialized_code_object, size_t* serialized_code_object_size) {
|
||||
return HsaApiTable->hsa_code_object_serialize_fn(
|
||||
return coreApiTable->hsa_code_object_serialize_fn(
|
||||
code_object, alloc_callback, callback_data, options,
|
||||
serialized_code_object, serialized_code_object_size);
|
||||
}
|
||||
@@ -470,33 +479,33 @@ hsa_status_t HSA_API
|
||||
size_t serialized_code_object_size,
|
||||
const char* options,
|
||||
hsa_code_object_t* code_object) {
|
||||
return HsaApiTable->hsa_code_object_deserialize_fn(
|
||||
return coreApiTable->hsa_code_object_deserialize_fn(
|
||||
serialized_code_object, serialized_code_object_size, options,
|
||||
code_object);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_code_object_destroy(hsa_code_object_t code_object) {
|
||||
return HsaApiTable->hsa_code_object_destroy_fn(code_object);
|
||||
return coreApiTable->hsa_code_object_destroy_fn(code_object);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_code_object_get_info(hsa_code_object_t code_object,
|
||||
hsa_code_object_info_t attribute,
|
||||
void* value) {
|
||||
return HsaApiTable->hsa_code_object_get_info_fn(code_object, attribute,
|
||||
return coreApiTable->hsa_code_object_get_info_fn(code_object, attribute,
|
||||
value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_code_object_get_symbol(hsa_code_object_t code_object,
|
||||
const char* symbol_name,
|
||||
hsa_code_symbol_t* symbol) {
|
||||
return HsaApiTable->hsa_code_object_get_symbol_fn(code_object, symbol_name,
|
||||
return coreApiTable->hsa_code_object_get_symbol_fn(code_object, symbol_name,
|
||||
symbol);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_code_symbol_get_info(hsa_code_symbol_t code_symbol,
|
||||
hsa_code_symbol_info_t attribute,
|
||||
void* value) {
|
||||
return HsaApiTable->hsa_code_symbol_get_info_fn(code_symbol, attribute,
|
||||
return coreApiTable->hsa_code_symbol_get_info_fn(code_symbol, attribute,
|
||||
value);
|
||||
}
|
||||
|
||||
@@ -505,7 +514,7 @@ hsa_status_t HSA_API hsa_code_object_iterate_symbols(
|
||||
hsa_status_t (*callback)(hsa_code_object_t code_object,
|
||||
hsa_code_symbol_t symbol, void* data),
|
||||
void* data) {
|
||||
return HsaApiTable->hsa_code_object_iterate_symbols_fn(code_object, callback,
|
||||
return coreApiTable->hsa_code_object_iterate_symbols_fn(code_object, callback,
|
||||
data);
|
||||
}
|
||||
|
||||
@@ -513,12 +522,12 @@ hsa_status_t HSA_API
|
||||
hsa_executable_create(hsa_profile_t profile,
|
||||
hsa_executable_state_t executable_state,
|
||||
const char* options, hsa_executable_t* executable) {
|
||||
return HsaApiTable->hsa_executable_create_fn(profile, executable_state,
|
||||
return coreApiTable->hsa_executable_create_fn(profile, executable_state,
|
||||
options, executable);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_executable_destroy(hsa_executable_t executable) {
|
||||
return HsaApiTable->hsa_executable_destroy_fn(executable);
|
||||
return coreApiTable->hsa_executable_destroy_fn(executable);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
@@ -526,26 +535,26 @@ hsa_status_t HSA_API
|
||||
hsa_agent_t agent,
|
||||
hsa_code_object_t code_object,
|
||||
const char* options) {
|
||||
return HsaApiTable->hsa_executable_load_code_object_fn(executable, agent,
|
||||
return coreApiTable->hsa_executable_load_code_object_fn(executable, agent,
|
||||
code_object, options);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_executable_freeze(hsa_executable_t executable, const char* options) {
|
||||
return HsaApiTable->hsa_executable_freeze_fn(executable, options);
|
||||
return coreApiTable->hsa_executable_freeze_fn(executable, options);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_executable_get_info(hsa_executable_t executable,
|
||||
hsa_executable_info_t attribute,
|
||||
void* value) {
|
||||
return HsaApiTable->hsa_executable_get_info_fn(executable, attribute, value);
|
||||
return coreApiTable->hsa_executable_get_info_fn(executable, attribute, value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_executable_global_variable_define(hsa_executable_t executable,
|
||||
const char* variable_name,
|
||||
void* address) {
|
||||
return HsaApiTable->hsa_executable_global_variable_define_fn(
|
||||
return coreApiTable->hsa_executable_global_variable_define_fn(
|
||||
executable, variable_name, address);
|
||||
}
|
||||
|
||||
@@ -554,7 +563,7 @@ hsa_status_t HSA_API
|
||||
hsa_agent_t agent,
|
||||
const char* variable_name,
|
||||
void* address) {
|
||||
return HsaApiTable->hsa_executable_agent_global_variable_define_fn(
|
||||
return coreApiTable->hsa_executable_agent_global_variable_define_fn(
|
||||
executable, agent, variable_name, address);
|
||||
}
|
||||
|
||||
@@ -563,13 +572,13 @@ hsa_status_t HSA_API
|
||||
hsa_agent_t agent,
|
||||
const char* variable_name,
|
||||
void* address) {
|
||||
return HsaApiTable->hsa_executable_readonly_variable_define_fn(
|
||||
return coreApiTable->hsa_executable_readonly_variable_define_fn(
|
||||
executable, agent, variable_name, address);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_executable_validate(hsa_executable_t executable, uint32_t* result) {
|
||||
return HsaApiTable->hsa_executable_validate_fn(executable, result);
|
||||
return coreApiTable->hsa_executable_validate_fn(executable, result);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
@@ -577,7 +586,7 @@ hsa_status_t HSA_API
|
||||
const char* module_name, const char* symbol_name,
|
||||
hsa_agent_t agent, int32_t call_convention,
|
||||
hsa_executable_symbol_t* symbol) {
|
||||
return HsaApiTable->hsa_executable_get_symbol_fn(
|
||||
return coreApiTable->hsa_executable_get_symbol_fn(
|
||||
executable, module_name, symbol_name, agent, call_convention, symbol);
|
||||
}
|
||||
|
||||
@@ -585,7 +594,7 @@ hsa_status_t HSA_API
|
||||
hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,
|
||||
hsa_executable_symbol_info_t attribute,
|
||||
void* value) {
|
||||
return HsaApiTable->hsa_executable_symbol_get_info_fn(executable_symbol,
|
||||
return coreApiTable->hsa_executable_symbol_get_info_fn(executable_symbol,
|
||||
attribute, value);
|
||||
}
|
||||
|
||||
@@ -594,11 +603,227 @@ hsa_status_t HSA_API hsa_executable_iterate_symbols(
|
||||
hsa_status_t (*callback)(hsa_executable_t executable,
|
||||
hsa_executable_symbol_t symbol, void* data),
|
||||
void* data) {
|
||||
return HsaApiTable->hsa_executable_iterate_symbols_fn(executable, callback,
|
||||
return coreApiTable->hsa_executable_iterate_symbols_fn(executable, callback,
|
||||
data);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_string(hsa_status_t status, const char** status_string) {
|
||||
return HsaApiTable->hsa_status_string_fn(status, status_string);
|
||||
return coreApiTable->hsa_status_string_fn(status, status_string);
|
||||
}
|
||||
|
||||
/*
|
||||
* Following set of functions are bundled as AMD Extension Apis
|
||||
*/
|
||||
|
||||
// Pass through stub functions
|
||||
hsa_status_t HSA_API hsa_amd_coherency_get_type(hsa_agent_t agent,
|
||||
hsa_amd_coherency_type_t* type) {
|
||||
return amdExtTable->hsa_amd_coherency_get_type_fn(agent, type);
|
||||
}
|
||||
|
||||
// Pass through stub functions
|
||||
hsa_status_t HSA_API hsa_amd_coherency_set_type(hsa_agent_t agent,
|
||||
hsa_amd_coherency_type_t type) {
|
||||
return amdExtTable->hsa_amd_coherency_set_type_fn(agent, type);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable) {
|
||||
return amdExtTable->hsa_amd_profiling_set_profiler_enabled_fn(
|
||||
queue, enable);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_profiling_async_copy_enable(bool enable) {
|
||||
return amdExtTable->hsa_amd_profiling_async_copy_enable_fn(enable);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_profiling_get_dispatch_time(
|
||||
hsa_agent_t agent, hsa_signal_t signal,
|
||||
hsa_amd_profiling_dispatch_time_t* time) {
|
||||
return amdExtTable->hsa_amd_profiling_get_dispatch_time_fn(
|
||||
agent, signal, time);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_profiling_get_async_copy_time(
|
||||
hsa_signal_t hsa_signal, hsa_amd_profiling_async_copy_time_t* time) {
|
||||
return amdExtTable->hsa_amd_profiling_get_async_copy_time_fn(hsa_signal, time);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent,
|
||||
uint64_t agent_tick,
|
||||
uint64_t* system_tick) {
|
||||
return amdExtTable->hsa_amd_profiling_convert_tick_to_system_domain_fn(
|
||||
agent, agent_tick, system_tick);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_signal_async_handler(hsa_signal_t signal,
|
||||
hsa_signal_condition_t cond,
|
||||
hsa_signal_value_t value,
|
||||
hsa_amd_signal_handler handler, void* arg) {
|
||||
return amdExtTable->hsa_amd_signal_async_handler_fn(
|
||||
signal, cond, value, handler, arg);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_async_function(void (*callback)(void* arg), void* arg) {
|
||||
return amdExtTable->hsa_amd_async_function_fn(callback, arg);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
uint32_t HSA_API
|
||||
hsa_amd_signal_wait_any(uint32_t signal_count, hsa_signal_t* signals,
|
||||
hsa_signal_condition_t* conds,
|
||||
hsa_signal_value_t* values, uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_hint,
|
||||
hsa_signal_value_t* satisfying_value) {
|
||||
return amdExtTable->hsa_amd_signal_wait_any_fn(
|
||||
signal_count, signals,
|
||||
conds, values, timeout_hint,
|
||||
wait_hint, satisfying_value);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_queue_cu_set_mask(const hsa_queue_t* queue,
|
||||
uint32_t num_cu_mask_count,
|
||||
const uint32_t* cu_mask) {
|
||||
return amdExtTable->hsa_amd_queue_cu_set_mask_fn(
|
||||
queue, num_cu_mask_count, cu_mask);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
|
||||
hsa_amd_memory_pool_info_t attribute,
|
||||
void* value) {
|
||||
return amdExtTable->hsa_amd_memory_pool_get_info_fn(
|
||||
memory_pool, attribute, value);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_agent_iterate_memory_pools(
|
||||
hsa_agent_t agent,
|
||||
hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void* data),
|
||||
void* data) {
|
||||
return amdExtTable->hsa_amd_agent_iterate_memory_pools_fn(
|
||||
agent, callback, data);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size,
|
||||
uint32_t flags, void** ptr) {
|
||||
return amdExtTable->hsa_amd_memory_pool_allocate_fn(
|
||||
memory_pool, size, flags, ptr);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_pool_free(void* ptr) {
|
||||
return amdExtTable->hsa_amd_memory_pool_free_fn(ptr);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent, const void* src,
|
||||
hsa_agent_t src_agent, size_t size,
|
||||
uint32_t num_dep_signals,
|
||||
const hsa_signal_t* dep_signals,
|
||||
hsa_signal_t completion_signal) {
|
||||
return amdExtTable->hsa_amd_memory_async_copy_fn(
|
||||
dst, dst_agent, src, src_agent, size,
|
||||
num_dep_signals, dep_signals, completion_signal);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_agent_memory_pool_get_info(
|
||||
hsa_agent_t agent, hsa_amd_memory_pool_t memory_pool,
|
||||
hsa_amd_agent_memory_pool_info_t attribute, void* value) {
|
||||
return amdExtTable->hsa_amd_agent_memory_pool_get_info_fn(
|
||||
agent, memory_pool, attribute, value);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents,
|
||||
const uint32_t* flags, const void* ptr) {
|
||||
return amdExtTable->hsa_amd_agents_allow_access_fn(
|
||||
num_agents, agents, flags, ptr);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_pool,
|
||||
hsa_amd_memory_pool_t dst_memory_pool,
|
||||
bool* result) {
|
||||
return amdExtTable->hsa_amd_memory_pool_can_migrate_fn(
|
||||
src_memory_pool, dst_memory_pool, result);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_migrate(const void* ptr,
|
||||
hsa_amd_memory_pool_t memory_pool,
|
||||
uint32_t flags) {
|
||||
return amdExtTable->hsa_amd_memory_migrate_fn(
|
||||
ptr, memory_pool, flags);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_lock(void* host_ptr, size_t size,
|
||||
hsa_agent_t* agents, int num_agent,
|
||||
void** agent_ptr) {
|
||||
return amdExtTable->hsa_amd_memory_lock_fn(
|
||||
host_ptr, size, agents, num_agent, agent_ptr);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_unlock(void* host_ptr) {
|
||||
return amdExtTable->hsa_amd_memory_unlock_fn(host_ptr);
|
||||
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) {
|
||||
return amdExtTable->hsa_amd_memory_fill_fn(ptr, value, count);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_interop_map_buffer(uint32_t num_agents,
|
||||
hsa_agent_t* agents,
|
||||
int interop_handle,
|
||||
uint32_t flags,
|
||||
size_t* size,
|
||||
void** ptr,
|
||||
size_t* metadata_size,
|
||||
const void** metadata) {
|
||||
return amdExtTable->hsa_amd_interop_map_buffer_fn(
|
||||
num_agents, agents, interop_handle,
|
||||
flags, size, ptr, metadata_size, metadata);
|
||||
}
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_interop_unmap_buffer(void* ptr) {
|
||||
return amdExtTable->hsa_amd_interop_unmap_buffer_fn(ptr);
|
||||
}
|
||||
|
||||
// Use the function pointer from local instance Image Extension
|
||||
hsa_status_t HSA_API hsa_amd_image_create(
|
||||
hsa_agent_t agent,
|
||||
const hsa_ext_image_descriptor_t *image_descriptor,
|
||||
const hsa_amd_image_descriptor_t *image_layout,
|
||||
const void *image_data,
|
||||
hsa_access_permission_t access_permission,
|
||||
hsa_ext_image_t *image) {
|
||||
return amdExtTable->hsa_amd_image_create_fn(agent, image_descriptor,
|
||||
image_layout, image_data, access_permission, image);
|
||||
}
|
||||
|
||||
|
||||
@@ -47,17 +47,28 @@
|
||||
#include "core/inc/hsa_internal.h"
|
||||
|
||||
namespace core {
|
||||
struct ApiTable {
|
||||
::ApiTable table;
|
||||
ExtTable extension_backup;
|
||||
struct HsaApiTable {
|
||||
|
||||
ApiTable();
|
||||
void Reset();
|
||||
void LinkExts(ExtTable* ptr);
|
||||
};
|
||||
static const uint32_t HSA_EXT_FINALIZER_API_TABLE_ID = 0;
|
||||
static const uint32_t HSA_EXT_IMAGE_API_TABLE_ID = 1;
|
||||
|
||||
extern ApiTable hsa_api_table_;
|
||||
extern ApiTable hsa_internal_api_table_;
|
||||
::HsaApiTable hsa_api;
|
||||
::CoreApiTable core_api;
|
||||
::AmdExtTable amd_ext_api;
|
||||
::FinalizerExtTable finalizer_api;
|
||||
::ImageExtTable image_api;
|
||||
|
||||
HsaApiTable();
|
||||
void Init();
|
||||
void UpdateCore();
|
||||
void UpdateAmdExts();
|
||||
void CloneExts(void* ptr, uint32_t table_id);
|
||||
void LinkExts(void* ptr, uint32_t table_id);
|
||||
void Reset();
|
||||
};
|
||||
|
||||
extern HsaApiTable hsa_api_table_;
|
||||
extern HsaApiTable hsa_internal_api_table_;
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
Ausführbare Datei
+186
@@ -0,0 +1,186 @@
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// The University of Illinois/NCSA
|
||||
// Open Source License (NCSA)
|
||||
//
|
||||
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
|
||||
//
|
||||
// Developed by:
|
||||
//
|
||||
// AMD Research and AMD HSA Software Development
|
||||
//
|
||||
// Advanced Micro Devices, Inc.
|
||||
//
|
||||
// www.amd.com
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to
|
||||
// deal with the Software without restriction, including without limitation
|
||||
// the rights to use, copy, modify, merge, publish, distribute, sublicense,
|
||||
// and/or sell copies of the Software, and to permit persons to whom the
|
||||
// Software is furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// - Redistributions of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimers.
|
||||
// - Redistributions in binary form must reproduce the above copyright
|
||||
// notice, this list of conditions and the following disclaimers in
|
||||
// the documentation and/or other materials provided with the distribution.
|
||||
// - Neither the names of Advanced Micro Devices, Inc,
|
||||
// nor the names of its contributors may be used to endorse or promote
|
||||
// products derived from this Software without specific prior written
|
||||
// permission.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
|
||||
// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
|
||||
// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
|
||||
// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
|
||||
// DEALINGS WITH THE SOFTWARE.
|
||||
//
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
// HSA AMD extension.
|
||||
|
||||
#ifndef HSA_RUNTIME_CORE_INC_EXT_AMD_H_
|
||||
#define HSA_RUNTIME_CORE_INC_EXT_AMD_H_
|
||||
|
||||
#include "hsa.h"
|
||||
#include "hsa_ext_image.h"
|
||||
#include "hsa_ext_amd.h"
|
||||
|
||||
// Wrap internal implementation inside AMD namespace
|
||||
namespace AMD {
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_coherency_get_type(hsa_agent_t agent,
|
||||
hsa_amd_coherency_type_t* type);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_coherency_set_type(hsa_agent_t agent,
|
||||
hsa_amd_coherency_type_t type);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_profiling_async_copy_enable(bool enable);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_profiling_get_dispatch_time(
|
||||
hsa_agent_t agent, hsa_signal_t signal,
|
||||
hsa_amd_profiling_dispatch_time_t* time);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_profiling_get_async_copy_time(
|
||||
hsa_signal_t signal, hsa_amd_profiling_async_copy_time_t* time);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent,
|
||||
uint64_t agent_tick,
|
||||
uint64_t* system_tick);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_signal_async_handler(hsa_signal_t signal,
|
||||
hsa_signal_condition_t cond,
|
||||
hsa_signal_value_t value,
|
||||
hsa_amd_signal_handler handler, void* arg);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_async_function(void (*callback)(void* arg), void* arg);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
uint32_t HSA_API
|
||||
hsa_amd_signal_wait_any(uint32_t signal_count, hsa_signal_t* signals,
|
||||
hsa_signal_condition_t* conds,
|
||||
hsa_signal_value_t* values, uint64_t timeout_hint,
|
||||
hsa_wait_state_t wait_hint,
|
||||
hsa_signal_value_t* satisfying_value);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_queue_cu_set_mask(const hsa_queue_t* queue,
|
||||
uint32_t num_cu_mask_count,
|
||||
const uint32_t* cu_mask);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
|
||||
hsa_amd_memory_pool_info_t attribute,
|
||||
void* value);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_agent_iterate_memory_pools(
|
||||
hsa_agent_t agent,
|
||||
hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void* data),
|
||||
void* data);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size,
|
||||
uint32_t flags, void** ptr);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_pool_free(void* ptr);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent, const void* src,
|
||||
hsa_agent_t src_agent, size_t size,
|
||||
uint32_t num_dep_signals,
|
||||
const hsa_signal_t* dep_signals,
|
||||
hsa_signal_t completion_signal);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_agent_memory_pool_get_info(
|
||||
hsa_agent_t agent, hsa_amd_memory_pool_t memory_pool,
|
||||
hsa_amd_agent_memory_pool_info_t attribute, void* value);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents,
|
||||
const uint32_t* flags, const void* ptr);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_pool,
|
||||
hsa_amd_memory_pool_t dst_memory_pool,
|
||||
bool* result);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_migrate(const void* ptr,
|
||||
hsa_amd_memory_pool_t memory_pool,
|
||||
uint32_t flags);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_lock(void* host_ptr, size_t size,
|
||||
hsa_agent_t* agents, int num_agent,
|
||||
void** agent_ptr);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_memory_unlock(void* host_ptr);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API
|
||||
hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_interop_map_buffer(uint32_t num_agents,
|
||||
hsa_agent_t* agents,
|
||||
int interop_handle,
|
||||
uint32_t flags,
|
||||
size_t* size,
|
||||
void** ptr,
|
||||
size_t* metadata_size,
|
||||
const void** metadata);
|
||||
|
||||
// Mirrors Amd Extension Apis
|
||||
hsa_status_t HSA_API hsa_amd_interop_unmap_buffer(void* ptr);
|
||||
|
||||
} // end of AMD namespace
|
||||
|
||||
#endif // header guard
|
||||
@@ -52,27 +52,43 @@
|
||||
#include "core/util/utils.h"
|
||||
|
||||
namespace core {
|
||||
struct ExtTableInternal : public ExtTable {
|
||||
struct ImageExtTableInternal : public ImageExtTable {
|
||||
decltype(::hsa_amd_image_get_info_max_dim)* hsa_amd_image_get_info_max_dim_fn;
|
||||
decltype(::hsa_amd_image_create)* hsa_amd_image_create_fn;
|
||||
};
|
||||
|
||||
class ExtensionEntryPoints {
|
||||
public:
|
||||
ExtTableInternal table;
|
||||
|
||||
// Table of function pointers for Hsa Extension Image
|
||||
ImageExtTableInternal image_api;
|
||||
|
||||
// Table of function pointers for Hsa Extension Finalizer
|
||||
FinalizerExtTable finalizer_api;
|
||||
|
||||
ExtensionEntryPoints();
|
||||
|
||||
bool Load(std::string library_name);
|
||||
bool LoadFinalizer(std::string library_name);
|
||||
bool LoadImage(std::string library_name);
|
||||
void Unload();
|
||||
|
||||
private:
|
||||
typedef void (*Load_t)(const ::ApiTable* table);
|
||||
typedef void (*Load_t)(const ::HsaApiTable* table);
|
||||
typedef void (*Unload_t)();
|
||||
|
||||
std::vector<os::LibHandle> libs_;
|
||||
|
||||
void InitTable();
|
||||
// Initialize table for HSA Finalizer Extension Api's
|
||||
void InitFinalizerExtTable();
|
||||
|
||||
// Initialize table for HSA Image Extension Api's
|
||||
void InitImageExtTable();
|
||||
|
||||
// Initialize Amd Ext table for Api related to Images
|
||||
void InitAmdExtTable();
|
||||
|
||||
// Update Amd Ext table for Api related to Images
|
||||
void UpdateAmdExtTable(void *func_ptr);
|
||||
|
||||
DISALLOW_COPY_AND_ASSIGN(ExtensionEntryPoints);
|
||||
};
|
||||
}
|
||||
|
||||
@@ -42,6 +42,6 @@
|
||||
|
||||
#include "hsa_api_trace.h"
|
||||
|
||||
void hsa_table_interface_init(const ApiTable* table);
|
||||
void hsa_table_interface_init(const HsaApiTable* apiTable);
|
||||
|
||||
const ApiTable* hsa_table_interface_get_table();
|
||||
const HsaApiTable* hsa_table_interface_get_table();
|
||||
|
||||
@@ -64,6 +64,7 @@
|
||||
#include "core/util/utils.h"
|
||||
#include "core/inc/registers.h"
|
||||
#include "core/inc/interrupt_signal.h"
|
||||
#include "core/inc/hsa_ext_amd_impl.h"
|
||||
|
||||
namespace amd {
|
||||
// Queue::amd_queue_ is cache-aligned for performance.
|
||||
@@ -249,7 +250,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id,
|
||||
}
|
||||
auto signal = new core::InterruptSignal(0, queue_event_);
|
||||
amd_queue_.queue_inactive_signal = core::InterruptSignal::Convert(signal);
|
||||
if (hsa_amd_signal_async_handler(
|
||||
if (AMD::hsa_amd_signal_async_handler(
|
||||
amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE, 0,
|
||||
DynamicScratchHandler, this) != HSA_STATUS_SUCCESS)
|
||||
return;
|
||||
|
||||
@@ -64,6 +64,8 @@
|
||||
// Size of scratch (private) segment pre-allocated per thread, in bytes.
|
||||
#define DEFAULT_SCRATCH_BYTES_PER_THREAD 2048
|
||||
|
||||
extern core::HsaApiTable hsa_internal_api_table_;
|
||||
|
||||
namespace amd {
|
||||
GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props)
|
||||
: GpuAgentInt(node),
|
||||
@@ -729,11 +731,11 @@ hsa_status_t GpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
|
||||
case HSA_AGENT_INFO_EXTENSIONS:
|
||||
memset(value, 0, sizeof(uint8_t) * 128);
|
||||
|
||||
if (extensions.table.hsa_ext_program_finalize_fn != NULL) {
|
||||
if (core::hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
|
||||
*((uint8_t*)value) = 1 << HSA_EXTENSION_FINALIZER;
|
||||
}
|
||||
|
||||
if (extensions.table.hsa_ext_image_create_fn != NULL) {
|
||||
if (core::hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
|
||||
*((uint8_t*)value) |= 1 << HSA_EXTENSION_IMAGES;
|
||||
}
|
||||
|
||||
|
||||
@@ -207,64 +207,67 @@ hsa_status_t
|
||||
hsa_status_t status = hsa_system_extension_supported(
|
||||
extension, version_major, version_minor, &supported);
|
||||
|
||||
if (HSA_STATUS_SUCCESS != status) {
|
||||
if ((HSA_STATUS_SUCCESS != status) ||
|
||||
(supported == false)) {
|
||||
return status;
|
||||
}
|
||||
|
||||
if (supported) {
|
||||
if (extension == HSA_EXTENSION_IMAGES) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ext_images_1_00_pfn_t* ext_table =
|
||||
reinterpret_cast<hsa_ext_images_1_00_pfn_t*>(table);
|
||||
ext_table->hsa_ext_image_clear = hsa_ext_image_clear;
|
||||
ext_table->hsa_ext_image_copy = hsa_ext_image_copy;
|
||||
ext_table->hsa_ext_image_create = hsa_ext_image_create;
|
||||
ext_table->hsa_ext_image_data_get_info = hsa_ext_image_data_get_info;
|
||||
ext_table->hsa_ext_image_destroy = hsa_ext_image_destroy;
|
||||
ext_table->hsa_ext_image_export = hsa_ext_image_export;
|
||||
ext_table->hsa_ext_image_get_capability = hsa_ext_image_get_capability;
|
||||
ext_table->hsa_ext_image_import = hsa_ext_image_import;
|
||||
ext_table->hsa_ext_sampler_create = hsa_ext_sampler_create;
|
||||
ext_table->hsa_ext_sampler_destroy = hsa_ext_sampler_destroy;
|
||||
if (extension == HSA_EXTENSION_IMAGES) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ext_images_1_00_pfn_t* ext_table =
|
||||
reinterpret_cast<hsa_ext_images_1_00_pfn_t*>(table);
|
||||
ext_table->hsa_ext_image_clear = hsa_ext_image_clear;
|
||||
ext_table->hsa_ext_image_copy = hsa_ext_image_copy;
|
||||
ext_table->hsa_ext_image_create = hsa_ext_image_create;
|
||||
ext_table->hsa_ext_image_data_get_info = hsa_ext_image_data_get_info;
|
||||
ext_table->hsa_ext_image_destroy = hsa_ext_image_destroy;
|
||||
ext_table->hsa_ext_image_export = hsa_ext_image_export;
|
||||
ext_table->hsa_ext_image_get_capability = hsa_ext_image_get_capability;
|
||||
ext_table->hsa_ext_image_import = hsa_ext_image_import;
|
||||
ext_table->hsa_ext_sampler_create = hsa_ext_sampler_create;
|
||||
ext_table->hsa_ext_sampler_destroy = hsa_ext_sampler_destroy;
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
} else if (extension == HSA_EXTENSION_FINALIZER) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ext_finalizer_1_00_pfn_s* ext_table =
|
||||
reinterpret_cast<hsa_ext_finalizer_1_00_pfn_s*>(table);
|
||||
ext_table->hsa_ext_program_add_module = hsa_ext_program_add_module;
|
||||
ext_table->hsa_ext_program_create = hsa_ext_program_create;
|
||||
ext_table->hsa_ext_program_destroy = hsa_ext_program_destroy;
|
||||
ext_table->hsa_ext_program_finalize = hsa_ext_program_finalize;
|
||||
ext_table->hsa_ext_program_get_info = hsa_ext_program_get_info;
|
||||
ext_table->hsa_ext_program_iterate_modules =
|
||||
hsa_ext_program_iterate_modules;
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
} else if (extension == HSA_EXTENSION_AMD_LOADED_CODE_OBJECT) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ven_amd_loaded_code_object_1_00_pfn_t* ext_table =
|
||||
reinterpret_cast<hsa_ven_amd_loaded_code_object_1_00_pfn_t*>(table);
|
||||
ext_table->hsa_ven_amd_loaded_code_object_query_host_address =
|
||||
hsa_ven_amd_loaded_code_object_query_host_address;
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
} else if (extension == HSA_EXTENSION_AMD_LOADER) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ven_amd_loader_1_00_pfn_t* ext_table =
|
||||
reinterpret_cast<hsa_ven_amd_loader_1_00_pfn_t*>(table);
|
||||
ext_table->hsa_ven_amd_loader_query_segment_descriptors =
|
||||
hsa_ven_amd_loader_query_segment_descriptors;
|
||||
ext_table->hsa_ven_amd_loader_query_host_address =
|
||||
hsa_ven_amd_loader_query_host_address;
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
} else {
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
if (extension == HSA_EXTENSION_FINALIZER) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ext_finalizer_1_00_pfn_s* ext_table =
|
||||
reinterpret_cast<hsa_ext_finalizer_1_00_pfn_s*>(table);
|
||||
ext_table->hsa_ext_program_add_module = hsa_ext_program_add_module;
|
||||
ext_table->hsa_ext_program_create = hsa_ext_program_create;
|
||||
ext_table->hsa_ext_program_destroy = hsa_ext_program_destroy;
|
||||
ext_table->hsa_ext_program_finalize = hsa_ext_program_finalize;
|
||||
ext_table->hsa_ext_program_get_info = hsa_ext_program_get_info;
|
||||
ext_table->hsa_ext_program_iterate_modules =
|
||||
hsa_ext_program_iterate_modules;
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
if (extension == HSA_EXTENSION_AMD_LOADED_CODE_OBJECT) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ven_amd_loaded_code_object_1_00_pfn_t* ext_table =
|
||||
reinterpret_cast<hsa_ven_amd_loaded_code_object_1_00_pfn_t*>(table);
|
||||
ext_table->hsa_ven_amd_loaded_code_object_query_host_address =
|
||||
hsa_ven_amd_loaded_code_object_query_host_address;
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
if (extension == HSA_EXTENSION_AMD_LOADER) {
|
||||
// Currently there is only version 1.00.
|
||||
hsa_ven_amd_loader_1_00_pfn_t* ext_table =
|
||||
reinterpret_cast<hsa_ven_amd_loader_1_00_pfn_t*>(table);
|
||||
ext_table->hsa_ven_amd_loader_query_segment_descriptors =
|
||||
hsa_ven_amd_loader_query_segment_descriptors;
|
||||
ext_table->hsa_ven_amd_loader_query_host_address =
|
||||
hsa_ven_amd_loader_query_host_address;
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
|
||||
//---------------------------------------------------------------------------//
|
||||
|
||||
@@ -42,150 +42,258 @@
|
||||
|
||||
#include "core/inc/hsa_api_trace_int.h"
|
||||
#include "core/inc/runtime.h"
|
||||
#include "core/inc/hsa_ext_amd_impl.h"
|
||||
#include "core/inc/hsa_table_interface.h"
|
||||
|
||||
#include <iostream>
|
||||
|
||||
namespace core {
|
||||
|
||||
ApiTable hsa_api_table_;
|
||||
ApiTable hsa_internal_api_table_;
|
||||
HsaApiTable hsa_api_table_;
|
||||
HsaApiTable hsa_internal_api_table_;
|
||||
|
||||
ApiTable::ApiTable() {
|
||||
table.std_exts_ = NULL;
|
||||
Reset();
|
||||
HsaApiTable::HsaApiTable() {
|
||||
Init();
|
||||
}
|
||||
|
||||
void ApiTable::LinkExts(ExtTable* ptr) {
|
||||
assert(ptr != NULL && "Invalid extension table linked.");
|
||||
extension_backup = *ptr;
|
||||
table.std_exts_ = ptr;
|
||||
// Initialize member fields for Hsa Core and Amd Extension Api's
|
||||
// Member fields for Finalizer and Image extensions will be
|
||||
// updated as part of Hsa Runtime initialization.
|
||||
void HsaApiTable::Init() {
|
||||
|
||||
// Initialize Version of Api Table
|
||||
hsa_api.version.major_id = HSA_API_TABLE_MAJOR_VERSION;
|
||||
hsa_api.version.minor_id = sizeof(::HsaApiTable);
|
||||
hsa_api.version.step_id = HSA_API_TABLE_STEP_VERSION;
|
||||
|
||||
// Update Api table for Core and its major id
|
||||
UpdateCore();
|
||||
hsa_api.core_ = &core_api;
|
||||
|
||||
// Update Api table for Amd Extensions and its major id
|
||||
UpdateAmdExts();
|
||||
hsa_api.amd_ext_ = &amd_ext_api;
|
||||
|
||||
// Initialize Api tables for Finalizer and Image to NULL
|
||||
// Tables for Finalizer and Images are initialized as part
|
||||
// of Hsa Runtime initialization, including their major ids
|
||||
hsa_api.finalizer_ext_ = NULL;
|
||||
hsa_api.image_ext_ = NULL;
|
||||
}
|
||||
|
||||
void ApiTable::Reset() {
|
||||
table.hsa_init_fn = HSA::hsa_init;
|
||||
table.hsa_shut_down_fn = HSA::hsa_shut_down;
|
||||
table.hsa_system_get_info_fn = HSA::hsa_system_get_info;
|
||||
table.hsa_system_extension_supported_fn = HSA::hsa_system_extension_supported;
|
||||
table.hsa_system_get_extension_table_fn = HSA::hsa_system_get_extension_table;
|
||||
table.hsa_iterate_agents_fn = HSA::hsa_iterate_agents;
|
||||
table.hsa_agent_get_info_fn = HSA::hsa_agent_get_info;
|
||||
table.hsa_agent_get_exception_policies_fn =
|
||||
void HsaApiTable::Reset() {
|
||||
Init();
|
||||
}
|
||||
|
||||
void HsaApiTable::CloneExts(void* ext_table, uint32_t table_id) {
|
||||
|
||||
assert(ext_table != NULL && "Invalid extension table linked.");
|
||||
|
||||
// Update HSA Extension Finalizer Api table
|
||||
if (table_id == HSA_EXT_FINALIZER_API_TABLE_ID) {
|
||||
finalizer_api = (*(FinalizerExtTable *)ext_table);
|
||||
hsa_api.finalizer_ext_ = &finalizer_api;
|
||||
return;
|
||||
}
|
||||
|
||||
// Update HSA Extension Image Api table
|
||||
if (table_id == HSA_EXT_IMAGE_API_TABLE_ID) {
|
||||
image_api = (*(ImageExtTable *)ext_table);
|
||||
hsa_api.image_ext_ = &image_api;
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
void HsaApiTable::LinkExts(void* ext_table, uint32_t table_id) {
|
||||
|
||||
assert(ext_table != NULL && "Invalid extension table linked.");
|
||||
|
||||
// Update HSA Extension Finalizer Api table
|
||||
if (table_id == HSA_EXT_FINALIZER_API_TABLE_ID) {
|
||||
finalizer_api = (*(FinalizerExtTable *)ext_table);
|
||||
hsa_api.finalizer_ext_ = (FinalizerExtTable *)ext_table;
|
||||
return;
|
||||
}
|
||||
|
||||
// Update HSA Extension Image Api table
|
||||
if (table_id == HSA_EXT_IMAGE_API_TABLE_ID) {
|
||||
image_api = (*(ImageExtTable *)ext_table);
|
||||
hsa_api.image_ext_ = (ImageExtTable *)ext_table;
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
// Update Api table for Hsa Core Runtime
|
||||
void HsaApiTable::UpdateCore() {
|
||||
|
||||
// Initialize Version of Api Table
|
||||
core_api.version.major_id = HSA_CORE_API_TABLE_MAJOR_VERSION;
|
||||
core_api.version.minor_id = sizeof(::CoreApiTable);
|
||||
core_api.version.step_id = HSA_CORE_API_TABLE_STEP_VERSION;
|
||||
|
||||
// Initialize function pointers for Hsa Core Runtime Api's
|
||||
core_api.hsa_init_fn = HSA::hsa_init;
|
||||
core_api.hsa_shut_down_fn = HSA::hsa_shut_down;
|
||||
core_api.hsa_system_get_info_fn = HSA::hsa_system_get_info;
|
||||
core_api.hsa_system_extension_supported_fn = HSA::hsa_system_extension_supported;
|
||||
core_api.hsa_system_get_extension_table_fn = HSA::hsa_system_get_extension_table;
|
||||
core_api.hsa_iterate_agents_fn = HSA::hsa_iterate_agents;
|
||||
core_api.hsa_agent_get_info_fn = HSA::hsa_agent_get_info;
|
||||
core_api.hsa_agent_get_exception_policies_fn =
|
||||
HSA::hsa_agent_get_exception_policies;
|
||||
table.hsa_agent_extension_supported_fn = HSA::hsa_agent_extension_supported;
|
||||
table.hsa_queue_create_fn = HSA::hsa_queue_create;
|
||||
table.hsa_soft_queue_create_fn = HSA::hsa_soft_queue_create;
|
||||
table.hsa_queue_destroy_fn = HSA::hsa_queue_destroy;
|
||||
table.hsa_queue_inactivate_fn = HSA::hsa_queue_inactivate;
|
||||
table.hsa_queue_load_read_index_acquire_fn =
|
||||
core_api.hsa_agent_extension_supported_fn = HSA::hsa_agent_extension_supported;
|
||||
core_api.hsa_queue_create_fn = HSA::hsa_queue_create;
|
||||
core_api.hsa_soft_queue_create_fn = HSA::hsa_soft_queue_create;
|
||||
core_api.hsa_queue_destroy_fn = HSA::hsa_queue_destroy;
|
||||
core_api.hsa_queue_inactivate_fn = HSA::hsa_queue_inactivate;
|
||||
core_api.hsa_queue_load_read_index_acquire_fn =
|
||||
HSA::hsa_queue_load_read_index_acquire;
|
||||
table.hsa_queue_load_read_index_relaxed_fn =
|
||||
core_api.hsa_queue_load_read_index_relaxed_fn =
|
||||
HSA::hsa_queue_load_read_index_relaxed;
|
||||
table.hsa_queue_load_write_index_acquire_fn =
|
||||
core_api.hsa_queue_load_write_index_acquire_fn =
|
||||
HSA::hsa_queue_load_write_index_acquire;
|
||||
table.hsa_queue_load_write_index_relaxed_fn =
|
||||
core_api.hsa_queue_load_write_index_relaxed_fn =
|
||||
HSA::hsa_queue_load_write_index_relaxed;
|
||||
table.hsa_queue_store_write_index_relaxed_fn =
|
||||
core_api.hsa_queue_store_write_index_relaxed_fn =
|
||||
HSA::hsa_queue_store_write_index_relaxed;
|
||||
table.hsa_queue_store_write_index_release_fn =
|
||||
core_api.hsa_queue_store_write_index_release_fn =
|
||||
HSA::hsa_queue_store_write_index_release;
|
||||
table.hsa_queue_cas_write_index_acq_rel_fn =
|
||||
core_api.hsa_queue_cas_write_index_acq_rel_fn =
|
||||
HSA::hsa_queue_cas_write_index_acq_rel;
|
||||
table.hsa_queue_cas_write_index_acquire_fn =
|
||||
core_api.hsa_queue_cas_write_index_acquire_fn =
|
||||
HSA::hsa_queue_cas_write_index_acquire;
|
||||
table.hsa_queue_cas_write_index_relaxed_fn =
|
||||
core_api.hsa_queue_cas_write_index_relaxed_fn =
|
||||
HSA::hsa_queue_cas_write_index_relaxed;
|
||||
table.hsa_queue_cas_write_index_release_fn =
|
||||
core_api.hsa_queue_cas_write_index_release_fn =
|
||||
HSA::hsa_queue_cas_write_index_release;
|
||||
table.hsa_queue_add_write_index_acq_rel_fn =
|
||||
core_api.hsa_queue_add_write_index_acq_rel_fn =
|
||||
HSA::hsa_queue_add_write_index_acq_rel;
|
||||
table.hsa_queue_add_write_index_acquire_fn =
|
||||
core_api.hsa_queue_add_write_index_acquire_fn =
|
||||
HSA::hsa_queue_add_write_index_acquire;
|
||||
table.hsa_queue_add_write_index_relaxed_fn =
|
||||
core_api.hsa_queue_add_write_index_relaxed_fn =
|
||||
HSA::hsa_queue_add_write_index_relaxed;
|
||||
table.hsa_queue_add_write_index_release_fn =
|
||||
core_api.hsa_queue_add_write_index_release_fn =
|
||||
HSA::hsa_queue_add_write_index_release;
|
||||
table.hsa_queue_store_read_index_relaxed_fn =
|
||||
core_api.hsa_queue_store_read_index_relaxed_fn =
|
||||
HSA::hsa_queue_store_read_index_relaxed;
|
||||
table.hsa_queue_store_read_index_release_fn =
|
||||
core_api.hsa_queue_store_read_index_release_fn =
|
||||
HSA::hsa_queue_store_read_index_release;
|
||||
table.hsa_agent_iterate_regions_fn = HSA::hsa_agent_iterate_regions;
|
||||
table.hsa_region_get_info_fn = HSA::hsa_region_get_info;
|
||||
table.hsa_memory_register_fn = HSA::hsa_memory_register;
|
||||
table.hsa_memory_deregister_fn = HSA::hsa_memory_deregister;
|
||||
table.hsa_memory_allocate_fn = HSA::hsa_memory_allocate;
|
||||
table.hsa_memory_free_fn = HSA::hsa_memory_free;
|
||||
table.hsa_memory_copy_fn = HSA::hsa_memory_copy;
|
||||
table.hsa_memory_assign_agent_fn = HSA::hsa_memory_assign_agent;
|
||||
table.hsa_signal_create_fn = HSA::hsa_signal_create;
|
||||
table.hsa_signal_destroy_fn = HSA::hsa_signal_destroy;
|
||||
table.hsa_signal_load_relaxed_fn = HSA::hsa_signal_load_relaxed;
|
||||
table.hsa_signal_load_acquire_fn = HSA::hsa_signal_load_acquire;
|
||||
table.hsa_signal_store_relaxed_fn = HSA::hsa_signal_store_relaxed;
|
||||
table.hsa_signal_store_release_fn = HSA::hsa_signal_store_release;
|
||||
table.hsa_signal_wait_relaxed_fn = HSA::hsa_signal_wait_relaxed;
|
||||
table.hsa_signal_wait_acquire_fn = HSA::hsa_signal_wait_acquire;
|
||||
table.hsa_signal_and_relaxed_fn = HSA::hsa_signal_and_relaxed;
|
||||
table.hsa_signal_and_acquire_fn = HSA::hsa_signal_and_acquire;
|
||||
table.hsa_signal_and_release_fn = HSA::hsa_signal_and_release;
|
||||
table.hsa_signal_and_acq_rel_fn = HSA::hsa_signal_and_acq_rel;
|
||||
table.hsa_signal_or_relaxed_fn = HSA::hsa_signal_or_relaxed;
|
||||
table.hsa_signal_or_acquire_fn = HSA::hsa_signal_or_acquire;
|
||||
table.hsa_signal_or_release_fn = HSA::hsa_signal_or_release;
|
||||
table.hsa_signal_or_acq_rel_fn = HSA::hsa_signal_or_acq_rel;
|
||||
table.hsa_signal_xor_relaxed_fn = HSA::hsa_signal_xor_relaxed;
|
||||
table.hsa_signal_xor_acquire_fn = HSA::hsa_signal_xor_acquire;
|
||||
table.hsa_signal_xor_release_fn = HSA::hsa_signal_xor_release;
|
||||
table.hsa_signal_xor_acq_rel_fn = HSA::hsa_signal_xor_acq_rel;
|
||||
table.hsa_signal_exchange_relaxed_fn = HSA::hsa_signal_exchange_relaxed;
|
||||
table.hsa_signal_exchange_acquire_fn = HSA::hsa_signal_exchange_acquire;
|
||||
table.hsa_signal_exchange_release_fn = HSA::hsa_signal_exchange_release;
|
||||
table.hsa_signal_exchange_acq_rel_fn = HSA::hsa_signal_exchange_acq_rel;
|
||||
table.hsa_signal_add_relaxed_fn = HSA::hsa_signal_add_relaxed;
|
||||
table.hsa_signal_add_acquire_fn = HSA::hsa_signal_add_acquire;
|
||||
table.hsa_signal_add_release_fn = HSA::hsa_signal_add_release;
|
||||
table.hsa_signal_add_acq_rel_fn = HSA::hsa_signal_add_acq_rel;
|
||||
table.hsa_signal_subtract_relaxed_fn = HSA::hsa_signal_subtract_relaxed;
|
||||
table.hsa_signal_subtract_acquire_fn = HSA::hsa_signal_subtract_acquire;
|
||||
table.hsa_signal_subtract_release_fn = HSA::hsa_signal_subtract_release;
|
||||
table.hsa_signal_subtract_acq_rel_fn = HSA::hsa_signal_subtract_acq_rel;
|
||||
table.hsa_signal_cas_relaxed_fn = HSA::hsa_signal_cas_relaxed;
|
||||
table.hsa_signal_cas_acquire_fn = HSA::hsa_signal_cas_acquire;
|
||||
table.hsa_signal_cas_release_fn = HSA::hsa_signal_cas_release;
|
||||
table.hsa_signal_cas_acq_rel_fn = HSA::hsa_signal_cas_acq_rel;
|
||||
table.hsa_isa_from_name_fn = HSA::hsa_isa_from_name;
|
||||
table.hsa_isa_get_info_fn = HSA::hsa_isa_get_info;
|
||||
table.hsa_isa_compatible_fn = HSA::hsa_isa_compatible;
|
||||
table.hsa_code_object_serialize_fn = HSA::hsa_code_object_serialize;
|
||||
table.hsa_code_object_deserialize_fn = HSA::hsa_code_object_deserialize;
|
||||
table.hsa_code_object_destroy_fn = HSA::hsa_code_object_destroy;
|
||||
table.hsa_code_object_get_info_fn = HSA::hsa_code_object_get_info;
|
||||
table.hsa_code_object_get_symbol_fn = HSA::hsa_code_object_get_symbol;
|
||||
table.hsa_code_symbol_get_info_fn = HSA::hsa_code_symbol_get_info;
|
||||
table.hsa_code_object_iterate_symbols_fn =
|
||||
core_api.hsa_agent_iterate_regions_fn = HSA::hsa_agent_iterate_regions;
|
||||
core_api.hsa_region_get_info_fn = HSA::hsa_region_get_info;
|
||||
core_api.hsa_memory_register_fn = HSA::hsa_memory_register;
|
||||
core_api.hsa_memory_deregister_fn = HSA::hsa_memory_deregister;
|
||||
core_api.hsa_memory_allocate_fn = HSA::hsa_memory_allocate;
|
||||
core_api.hsa_memory_free_fn = HSA::hsa_memory_free;
|
||||
core_api.hsa_memory_copy_fn = HSA::hsa_memory_copy;
|
||||
core_api.hsa_memory_assign_agent_fn = HSA::hsa_memory_assign_agent;
|
||||
core_api.hsa_signal_create_fn = HSA::hsa_signal_create;
|
||||
core_api.hsa_signal_destroy_fn = HSA::hsa_signal_destroy;
|
||||
core_api.hsa_signal_load_relaxed_fn = HSA::hsa_signal_load_relaxed;
|
||||
core_api.hsa_signal_load_acquire_fn = HSA::hsa_signal_load_acquire;
|
||||
core_api.hsa_signal_store_relaxed_fn = HSA::hsa_signal_store_relaxed;
|
||||
core_api.hsa_signal_store_release_fn = HSA::hsa_signal_store_release;
|
||||
core_api.hsa_signal_wait_relaxed_fn = HSA::hsa_signal_wait_relaxed;
|
||||
core_api.hsa_signal_wait_acquire_fn = HSA::hsa_signal_wait_acquire;
|
||||
core_api.hsa_signal_and_relaxed_fn = HSA::hsa_signal_and_relaxed;
|
||||
core_api.hsa_signal_and_acquire_fn = HSA::hsa_signal_and_acquire;
|
||||
core_api.hsa_signal_and_release_fn = HSA::hsa_signal_and_release;
|
||||
core_api.hsa_signal_and_acq_rel_fn = HSA::hsa_signal_and_acq_rel;
|
||||
core_api.hsa_signal_or_relaxed_fn = HSA::hsa_signal_or_relaxed;
|
||||
core_api.hsa_signal_or_acquire_fn = HSA::hsa_signal_or_acquire;
|
||||
core_api.hsa_signal_or_release_fn = HSA::hsa_signal_or_release;
|
||||
core_api.hsa_signal_or_acq_rel_fn = HSA::hsa_signal_or_acq_rel;
|
||||
core_api.hsa_signal_xor_relaxed_fn = HSA::hsa_signal_xor_relaxed;
|
||||
core_api.hsa_signal_xor_acquire_fn = HSA::hsa_signal_xor_acquire;
|
||||
core_api.hsa_signal_xor_release_fn = HSA::hsa_signal_xor_release;
|
||||
core_api.hsa_signal_xor_acq_rel_fn = HSA::hsa_signal_xor_acq_rel;
|
||||
core_api.hsa_signal_exchange_relaxed_fn = HSA::hsa_signal_exchange_relaxed;
|
||||
core_api.hsa_signal_exchange_acquire_fn = HSA::hsa_signal_exchange_acquire;
|
||||
core_api.hsa_signal_exchange_release_fn = HSA::hsa_signal_exchange_release;
|
||||
core_api.hsa_signal_exchange_acq_rel_fn = HSA::hsa_signal_exchange_acq_rel;
|
||||
core_api.hsa_signal_add_relaxed_fn = HSA::hsa_signal_add_relaxed;
|
||||
core_api.hsa_signal_add_acquire_fn = HSA::hsa_signal_add_acquire;
|
||||
core_api.hsa_signal_add_release_fn = HSA::hsa_signal_add_release;
|
||||
core_api.hsa_signal_add_acq_rel_fn = HSA::hsa_signal_add_acq_rel;
|
||||
core_api.hsa_signal_subtract_relaxed_fn = HSA::hsa_signal_subtract_relaxed;
|
||||
core_api.hsa_signal_subtract_acquire_fn = HSA::hsa_signal_subtract_acquire;
|
||||
core_api.hsa_signal_subtract_release_fn = HSA::hsa_signal_subtract_release;
|
||||
core_api.hsa_signal_subtract_acq_rel_fn = HSA::hsa_signal_subtract_acq_rel;
|
||||
core_api.hsa_signal_cas_relaxed_fn = HSA::hsa_signal_cas_relaxed;
|
||||
core_api.hsa_signal_cas_acquire_fn = HSA::hsa_signal_cas_acquire;
|
||||
core_api.hsa_signal_cas_release_fn = HSA::hsa_signal_cas_release;
|
||||
core_api.hsa_signal_cas_acq_rel_fn = HSA::hsa_signal_cas_acq_rel;
|
||||
core_api.hsa_isa_from_name_fn = HSA::hsa_isa_from_name;
|
||||
core_api.hsa_isa_get_info_fn = HSA::hsa_isa_get_info;
|
||||
core_api.hsa_isa_compatible_fn = HSA::hsa_isa_compatible;
|
||||
core_api.hsa_code_object_serialize_fn = HSA::hsa_code_object_serialize;
|
||||
core_api.hsa_code_object_deserialize_fn = HSA::hsa_code_object_deserialize;
|
||||
core_api.hsa_code_object_destroy_fn = HSA::hsa_code_object_destroy;
|
||||
core_api.hsa_code_object_get_info_fn = HSA::hsa_code_object_get_info;
|
||||
core_api.hsa_code_object_get_symbol_fn = HSA::hsa_code_object_get_symbol;
|
||||
core_api.hsa_code_symbol_get_info_fn = HSA::hsa_code_symbol_get_info;
|
||||
core_api.hsa_code_object_iterate_symbols_fn =
|
||||
HSA::hsa_code_object_iterate_symbols;
|
||||
table.hsa_executable_create_fn = HSA::hsa_executable_create;
|
||||
table.hsa_executable_destroy_fn = HSA::hsa_executable_destroy;
|
||||
table.hsa_executable_load_code_object_fn =
|
||||
core_api.hsa_executable_create_fn = HSA::hsa_executable_create;
|
||||
core_api.hsa_executable_destroy_fn = HSA::hsa_executable_destroy;
|
||||
core_api.hsa_executable_load_code_object_fn =
|
||||
HSA::hsa_executable_load_code_object;
|
||||
table.hsa_executable_freeze_fn = HSA::hsa_executable_freeze;
|
||||
table.hsa_executable_get_info_fn = HSA::hsa_executable_get_info;
|
||||
table.hsa_executable_global_variable_define_fn =
|
||||
core_api.hsa_executable_freeze_fn = HSA::hsa_executable_freeze;
|
||||
core_api.hsa_executable_get_info_fn = HSA::hsa_executable_get_info;
|
||||
core_api.hsa_executable_global_variable_define_fn =
|
||||
HSA::hsa_executable_global_variable_define;
|
||||
table.hsa_executable_agent_global_variable_define_fn =
|
||||
core_api.hsa_executable_agent_global_variable_define_fn =
|
||||
HSA::hsa_executable_agent_global_variable_define;
|
||||
table.hsa_executable_readonly_variable_define_fn =
|
||||
core_api.hsa_executable_readonly_variable_define_fn =
|
||||
HSA::hsa_executable_readonly_variable_define;
|
||||
table.hsa_executable_validate_fn = HSA::hsa_executable_validate;
|
||||
table.hsa_executable_get_symbol_fn = HSA::hsa_executable_get_symbol;
|
||||
table.hsa_executable_symbol_get_info_fn = HSA::hsa_executable_symbol_get_info;
|
||||
table.hsa_executable_iterate_symbols_fn = HSA::hsa_executable_iterate_symbols;
|
||||
table.hsa_status_string_fn = HSA::hsa_status_string;
|
||||
core_api.hsa_executable_validate_fn = HSA::hsa_executable_validate;
|
||||
core_api.hsa_executable_get_symbol_fn = HSA::hsa_executable_get_symbol;
|
||||
core_api.hsa_executable_symbol_get_info_fn = HSA::hsa_executable_symbol_get_info;
|
||||
core_api.hsa_executable_iterate_symbols_fn = HSA::hsa_executable_iterate_symbols;
|
||||
core_api.hsa_status_string_fn = HSA::hsa_status_string;
|
||||
}
|
||||
|
||||
if (table.std_exts_ != NULL) *table.std_exts_ = extension_backup;
|
||||
// Update Api table for Amd Extensions.
|
||||
// @note: Current implementation will initialize the
|
||||
// member variable hsa_amd_image_create_fn while loading
|
||||
// Image extension library
|
||||
void HsaApiTable::UpdateAmdExts() {
|
||||
|
||||
// Initialize Version of Api Table
|
||||
amd_ext_api.version.major_id = HSA_AMD_EXT_API_TABLE_MAJOR_VERSION;
|
||||
amd_ext_api.version.minor_id = sizeof(::AmdExtTable);
|
||||
amd_ext_api.version.step_id = HSA_AMD_EXT_API_TABLE_STEP_VERSION;
|
||||
|
||||
// Initialize function pointers for Amd Extension Api's
|
||||
amd_ext_api.hsa_amd_coherency_get_type_fn = AMD::hsa_amd_coherency_get_type;
|
||||
amd_ext_api.hsa_amd_coherency_set_type_fn = AMD::hsa_amd_coherency_set_type;
|
||||
amd_ext_api.hsa_amd_profiling_set_profiler_enabled_fn = AMD::hsa_amd_profiling_set_profiler_enabled;
|
||||
amd_ext_api.hsa_amd_profiling_async_copy_enable_fn = AMD::hsa_amd_profiling_async_copy_enable;
|
||||
amd_ext_api.hsa_amd_profiling_get_dispatch_time_fn = AMD::hsa_amd_profiling_get_dispatch_time;
|
||||
amd_ext_api.hsa_amd_profiling_get_async_copy_time_fn = AMD::hsa_amd_profiling_get_async_copy_time;
|
||||
amd_ext_api.hsa_amd_profiling_convert_tick_to_system_domain_fn = AMD::hsa_amd_profiling_convert_tick_to_system_domain;
|
||||
amd_ext_api.hsa_amd_signal_async_handler_fn = AMD::hsa_amd_signal_async_handler;
|
||||
amd_ext_api.hsa_amd_async_function_fn = AMD::hsa_amd_async_function;
|
||||
amd_ext_api.hsa_amd_signal_wait_any_fn = AMD::hsa_amd_signal_wait_any;
|
||||
amd_ext_api.hsa_amd_queue_cu_set_mask_fn = AMD::hsa_amd_queue_cu_set_mask;
|
||||
amd_ext_api.hsa_amd_memory_pool_get_info_fn = AMD::hsa_amd_memory_pool_get_info;
|
||||
amd_ext_api.hsa_amd_agent_iterate_memory_pools_fn = AMD::hsa_amd_agent_iterate_memory_pools;
|
||||
amd_ext_api.hsa_amd_memory_pool_allocate_fn = AMD::hsa_amd_memory_pool_allocate;
|
||||
amd_ext_api.hsa_amd_memory_pool_free_fn = AMD::hsa_amd_memory_pool_free;
|
||||
amd_ext_api.hsa_amd_memory_async_copy_fn = AMD::hsa_amd_memory_async_copy;
|
||||
amd_ext_api.hsa_amd_agent_memory_pool_get_info_fn = AMD::hsa_amd_agent_memory_pool_get_info;
|
||||
amd_ext_api.hsa_amd_agents_allow_access_fn = AMD::hsa_amd_agents_allow_access;
|
||||
amd_ext_api.hsa_amd_memory_pool_can_migrate_fn = AMD::hsa_amd_memory_pool_can_migrate;
|
||||
amd_ext_api.hsa_amd_memory_migrate_fn = AMD::hsa_amd_memory_migrate;
|
||||
amd_ext_api.hsa_amd_memory_lock_fn = AMD::hsa_amd_memory_lock;
|
||||
amd_ext_api.hsa_amd_memory_unlock_fn = AMD::hsa_amd_memory_unlock;
|
||||
amd_ext_api.hsa_amd_memory_fill_fn = AMD::hsa_amd_memory_fill;
|
||||
amd_ext_api.hsa_amd_interop_map_buffer_fn = AMD::hsa_amd_interop_map_buffer;
|
||||
amd_ext_api.hsa_amd_interop_unmap_buffer_fn = AMD::hsa_amd_interop_unmap_buffer;
|
||||
}
|
||||
|
||||
class Init {
|
||||
public:
|
||||
Init() { hsa_table_interface_init(&hsa_api_table_.table); }
|
||||
Init() { hsa_table_interface_init(&hsa_api_table_.hsa_api); }
|
||||
};
|
||||
static Init LinkAtLoad;
|
||||
}
|
||||
|
||||
@@ -109,7 +109,9 @@ static __forceinline bool IsValid(T* ptr) {
|
||||
return (ptr == NULL) ? NULL : ptr->IsValid();
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
namespace AMD {
|
||||
|
||||
hsa_status_t
|
||||
hsa_amd_coherency_get_type(hsa_agent_t agent_handle,
|
||||
hsa_amd_coherency_type_t* type) {
|
||||
IS_OPEN();
|
||||
@@ -132,7 +134,7 @@ hsa_status_t HSA_API
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_coherency_set_type(hsa_agent_t agent_handle,
|
||||
hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle,
|
||||
hsa_amd_coherency_type_t type) {
|
||||
IS_OPEN();
|
||||
|
||||
@@ -158,7 +160,7 @@ hsa_status_t HSA_API hsa_amd_coherency_set_type(hsa_agent_t agent_handle,
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) {
|
||||
IS_OPEN();
|
||||
|
||||
@@ -173,7 +175,7 @@ hsa_status_t HSA_API
|
||||
return core::Runtime::runtime_singleton_->FillMemory(ptr, value, count);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent_handle,
|
||||
const void* src, hsa_agent_t src_agent_handle,
|
||||
size_t size, uint32_t num_dep_signals,
|
||||
@@ -215,7 +217,7 @@ hsa_status_t HSA_API
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable) {
|
||||
IS_OPEN();
|
||||
|
||||
@@ -229,7 +231,7 @@ hsa_status_t HSA_API
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_profiling_async_copy_enable(bool enable) {
|
||||
hsa_status_t hsa_amd_profiling_async_copy_enable(bool enable) {
|
||||
IS_OPEN();
|
||||
|
||||
return core::Runtime::runtime_singleton_->IterateAgent(
|
||||
@@ -240,7 +242,7 @@ hsa_status_t HSA_API hsa_amd_profiling_async_copy_enable(bool enable) {
|
||||
reinterpret_cast<void*>(&enable));
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_profiling_get_dispatch_time(
|
||||
hsa_status_t hsa_amd_profiling_get_dispatch_time(
|
||||
hsa_agent_t agent_handle, hsa_signal_t hsa_signal,
|
||||
hsa_amd_profiling_dispatch_time_t* time) {
|
||||
IS_OPEN();
|
||||
@@ -267,7 +269,7 @@ hsa_status_t HSA_API hsa_amd_profiling_get_dispatch_time(
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_profiling_get_async_copy_time(
|
||||
hsa_status_t hsa_amd_profiling_get_async_copy_time(
|
||||
hsa_signal_t hsa_signal, hsa_amd_profiling_async_copy_time_t* time) {
|
||||
IS_OPEN();
|
||||
|
||||
@@ -295,7 +297,7 @@ hsa_status_t HSA_API hsa_amd_profiling_get_async_copy_time(
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent_handle,
|
||||
uint64_t agent_tick,
|
||||
uint64_t* system_tick) {
|
||||
@@ -318,7 +320,7 @@ hsa_status_t HSA_API
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
uint32_t HSA_API
|
||||
uint32_t
|
||||
hsa_amd_signal_wait_any(uint32_t signal_count, hsa_signal_t* hsa_signals,
|
||||
hsa_signal_condition_t* conds,
|
||||
hsa_signal_value_t* values, uint64_t timeout_hint,
|
||||
@@ -336,7 +338,7 @@ uint32_t HSA_API
|
||||
timeout_hint, wait_hint, satisfying_value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_signal_async_handler(hsa_signal_t hsa_signal,
|
||||
hsa_signal_condition_t cond,
|
||||
hsa_signal_value_t value,
|
||||
@@ -352,7 +354,7 @@ hsa_status_t HSA_API
|
||||
hsa_signal, cond, value, handler, arg);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_async_function(void (*callback)(void* arg), void* arg) {
|
||||
IS_OPEN();
|
||||
|
||||
@@ -363,7 +365,7 @@ hsa_status_t HSA_API
|
||||
arg);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_queue_cu_set_mask(const hsa_queue_t* queue,
|
||||
hsa_status_t hsa_amd_queue_cu_set_mask(const hsa_queue_t* queue,
|
||||
uint32_t num_cu_mask_count,
|
||||
const uint32_t* cu_mask) {
|
||||
IS_OPEN();
|
||||
@@ -374,7 +376,7 @@ hsa_status_t HSA_API hsa_amd_queue_cu_set_mask(const hsa_queue_t* queue,
|
||||
return cmd_queue->SetCUMasking(num_cu_mask_count, cu_mask);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_memory_lock(void* host_ptr, size_t size,
|
||||
hsa_status_t hsa_amd_memory_lock(void* host_ptr, size_t size,
|
||||
hsa_agent_t* agents, int num_agent,
|
||||
void** agent_ptr) {
|
||||
*agent_ptr = NULL;
|
||||
@@ -397,7 +399,7 @@ hsa_status_t HSA_API hsa_amd_memory_lock(void* host_ptr, size_t size,
|
||||
return system_region->Lock(num_agent, agents, host_ptr, size, agent_ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_memory_unlock(void* host_ptr) {
|
||||
hsa_status_t hsa_amd_memory_unlock(void* host_ptr) {
|
||||
IS_OPEN();
|
||||
|
||||
const amd::MemoryRegion* system_region =
|
||||
@@ -407,7 +409,7 @@ hsa_status_t HSA_API hsa_amd_memory_unlock(void* host_ptr) {
|
||||
return system_region->Unlock(host_ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
|
||||
hsa_amd_memory_pool_info_t attribute,
|
||||
void* value) {
|
||||
@@ -423,7 +425,7 @@ hsa_status_t HSA_API
|
||||
return mem_region->GetPoolInfo(attribute, value);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_agent_iterate_memory_pools(
|
||||
hsa_status_t hsa_amd_agent_iterate_memory_pools(
|
||||
hsa_agent_t agent_handle,
|
||||
hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void* data),
|
||||
void* data) {
|
||||
@@ -446,7 +448,7 @@ hsa_status_t HSA_API hsa_amd_agent_iterate_memory_pools(
|
||||
data);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size,
|
||||
uint32_t flags, void** ptr) {
|
||||
IS_OPEN();
|
||||
@@ -466,11 +468,11 @@ hsa_status_t HSA_API
|
||||
size, ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_memory_pool_free(void* ptr) {
|
||||
hsa_status_t hsa_amd_memory_pool_free(void* ptr) {
|
||||
return HSA::hsa_memory_free(ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents,
|
||||
const uint32_t* flags, const void* ptr) {
|
||||
IS_OPEN();
|
||||
@@ -483,7 +485,7 @@ hsa_status_t HSA_API
|
||||
ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API
|
||||
hsa_status_t
|
||||
hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_pool,
|
||||
hsa_amd_memory_pool_t dst_memory_pool,
|
||||
bool* result) {
|
||||
@@ -512,7 +514,7 @@ hsa_status_t HSA_API
|
||||
return src_mem_region->CanMigrate(*dst_mem_region, *result);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_memory_migrate(const void* ptr,
|
||||
hsa_status_t hsa_amd_memory_migrate(const void* ptr,
|
||||
hsa_amd_memory_pool_t memory_pool,
|
||||
uint32_t flags) {
|
||||
IS_OPEN();
|
||||
@@ -532,7 +534,7 @@ hsa_status_t HSA_API hsa_amd_memory_migrate(const void* ptr,
|
||||
return dst_mem_region->Migrate(flags, ptr);
|
||||
}
|
||||
|
||||
hsa_status_t HSA_API hsa_amd_agent_memory_pool_get_info(
|
||||
hsa_status_t hsa_amd_agent_memory_pool_get_info(
|
||||
hsa_agent_t agent_handle, hsa_amd_memory_pool_t memory_pool,
|
||||
hsa_amd_agent_memory_pool_info_t attribute, void* value) {
|
||||
IS_OPEN();
|
||||
@@ -593,3 +595,5 @@ hsa_status_t hsa_amd_interop_unmap_buffer(void* ptr) {
|
||||
if (ptr != NULL) core::Runtime::runtime_singleton_->InteropUnmap(ptr);
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
} // end of AMD namespace
|
||||
|
||||
+208
-134
@@ -161,28 +161,75 @@ static T0 hsa_ext_null(T1, T2, T3, T4, T5, T6, T7, T8, T9, T10, T11, T12, T13,
|
||||
T14, T15, T16, T17, T18, T19, T20) {
|
||||
return HSA_STATUS_ERROR_NOT_INITIALIZED;
|
||||
}
|
||||
template <class T0, class T1, class T2, class T3, class T4, class T5, class T6>
|
||||
static T0 hsa_amd_null(T1, T2, T3, T4, T5, T6) {
|
||||
return HSA_STATUS_ERROR_NOT_INITIALIZED;
|
||||
}
|
||||
|
||||
ExtensionEntryPoints::ExtensionEntryPoints() { InitTable(); }
|
||||
ExtensionEntryPoints::ExtensionEntryPoints() {
|
||||
InitFinalizerExtTable();
|
||||
InitImageExtTable();
|
||||
InitAmdExtTable();
|
||||
}
|
||||
|
||||
void ExtensionEntryPoints::InitTable() {
|
||||
table.hsa_ext_program_create_fn = hsa_ext_null;
|
||||
table.hsa_ext_program_destroy_fn = hsa_ext_null;
|
||||
table.hsa_ext_program_add_module_fn = hsa_ext_null;
|
||||
table.hsa_ext_program_iterate_modules_fn = hsa_ext_null;
|
||||
table.hsa_ext_program_get_info_fn = hsa_ext_null;
|
||||
table.hsa_ext_program_finalize_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_get_capability_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_data_get_info_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_create_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_import_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_export_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_copy_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_clear_fn = hsa_ext_null;
|
||||
table.hsa_ext_image_destroy_fn = hsa_ext_null;
|
||||
table.hsa_ext_sampler_create_fn = hsa_ext_null;
|
||||
table.hsa_ext_sampler_destroy_fn = hsa_ext_null;
|
||||
table.hsa_amd_image_get_info_max_dim_fn = hsa_ext_null;
|
||||
table.hsa_amd_image_create_fn = hsa_ext_null;
|
||||
// Initialize Finalizer function table to be NULLs
|
||||
void ExtensionEntryPoints::InitFinalizerExtTable() {
|
||||
|
||||
// Initialize Version of Api Table
|
||||
finalizer_api.version.major_id = 0x00;
|
||||
finalizer_api.version.minor_id = 0x00;
|
||||
finalizer_api.version.step_id = 0x00;
|
||||
|
||||
finalizer_api.hsa_ext_program_create_fn = hsa_ext_null;
|
||||
finalizer_api.hsa_ext_program_destroy_fn = hsa_ext_null;
|
||||
finalizer_api.hsa_ext_program_add_module_fn = hsa_ext_null;
|
||||
finalizer_api.hsa_ext_program_iterate_modules_fn = hsa_ext_null;
|
||||
finalizer_api.hsa_ext_program_get_info_fn = hsa_ext_null;
|
||||
finalizer_api.hsa_ext_program_finalize_fn = hsa_ext_null;
|
||||
}
|
||||
|
||||
// Initialize Image function table to be NULLs
|
||||
void ExtensionEntryPoints::InitImageExtTable() {
|
||||
|
||||
// Initialize Version of Api Table
|
||||
image_api.version.major_id = 0x00;
|
||||
image_api.version.minor_id = 0x00;
|
||||
image_api.version.step_id = 0x00;
|
||||
|
||||
image_api.hsa_ext_image_get_capability_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_image_data_get_info_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_image_create_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_image_import_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_image_export_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_image_copy_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_image_clear_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_image_destroy_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_sampler_create_fn = hsa_ext_null;
|
||||
image_api.hsa_ext_sampler_destroy_fn = hsa_ext_null;
|
||||
image_api.hsa_amd_image_get_info_max_dim_fn = hsa_ext_null;
|
||||
}
|
||||
|
||||
// Initialize Amd Ext table for Api related to Images
|
||||
void ExtensionEntryPoints::InitAmdExtTable() {
|
||||
hsa_api_table_.amd_ext_api.hsa_amd_image_create_fn = hsa_ext_null;
|
||||
hsa_internal_api_table_.amd_ext_api.hsa_amd_image_create_fn = hsa_ext_null;
|
||||
}
|
||||
|
||||
// Update Amd Ext table for Api related to Images.
|
||||
// @note: Interface should be updated when Amd Ext table
|
||||
// begins hosting Api's from other extension libraries
|
||||
void ExtensionEntryPoints::UpdateAmdExtTable(void *func_ptr) {
|
||||
|
||||
assert(hsa_api_table_.amd_ext_api.hsa_amd_image_create_fn ==
|
||||
(decltype(::hsa_amd_image_create)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
assert(hsa_internal_api_table_.amd_ext_api.hsa_amd_image_create_fn ==
|
||||
(decltype(::hsa_amd_image_create)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
hsa_api_table_.amd_ext_api.hsa_amd_image_create_fn =
|
||||
(decltype(::hsa_amd_image_create)*)func_ptr;
|
||||
hsa_internal_api_table_.amd_ext_api.hsa_amd_image_create_fn =
|
||||
(decltype(::hsa_amd_image_create)*)func_ptr;
|
||||
}
|
||||
|
||||
void ExtensionEntryPoints::Unload() {
|
||||
@@ -200,182 +247,217 @@ void ExtensionEntryPoints::Unload() {
|
||||
}
|
||||
}
|
||||
libs_.clear();
|
||||
InitTable();
|
||||
|
||||
InitFinalizerExtTable();
|
||||
InitImageExtTable();
|
||||
InitAmdExtTable();
|
||||
core::hsa_internal_api_table_.Reset();
|
||||
}
|
||||
|
||||
bool ExtensionEntryPoints::Load(std::string library_name) {
|
||||
bool ExtensionEntryPoints::LoadImage(std::string library_name) {
|
||||
os::LibHandle lib = os::LoadLib(library_name);
|
||||
if (lib == NULL) {
|
||||
return false;
|
||||
}
|
||||
libs_.push_back(lib);
|
||||
|
||||
|
||||
void* ptr;
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_create_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_program_create_fn ==
|
||||
(decltype(::hsa_ext_program_create)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_program_create_fn = (decltype(::hsa_ext_program_create)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_destroy_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_program_destroy_fn ==
|
||||
(decltype(::hsa_ext_program_destroy)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_program_destroy_fn =
|
||||
(decltype(::hsa_ext_program_destroy)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_add_module_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_program_add_module_fn ==
|
||||
(decltype(::hsa_ext_program_add_module)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_program_add_module_fn =
|
||||
(decltype(::hsa_ext_program_add_module)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_iterate_modules_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_program_iterate_modules_fn ==
|
||||
(decltype(::hsa_ext_program_iterate_modules)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_program_iterate_modules_fn =
|
||||
(decltype(::hsa_ext_program_iterate_modules)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_get_info_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_program_get_info_fn ==
|
||||
(decltype(::hsa_ext_program_get_info)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_program_get_info_fn =
|
||||
(decltype(::hsa_ext_program_get_info)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_finalize_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_program_finalize_fn ==
|
||||
(decltype(::hsa_ext_program_finalize)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_program_finalize_fn =
|
||||
(decltype(::hsa_ext_program_finalize)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_get_capability_impl");
|
||||
bool libIsImage = (ptr != NULL);
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_get_capability_fn ==
|
||||
assert(image_api.hsa_ext_image_get_capability_fn ==
|
||||
(decltype(::hsa_ext_image_get_capability)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_get_capability_fn =
|
||||
image_api.hsa_ext_image_get_capability_fn =
|
||||
(decltype(::hsa_ext_image_get_capability)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_data_get_info_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_data_get_info_fn ==
|
||||
assert(image_api.hsa_ext_image_data_get_info_fn ==
|
||||
(decltype(::hsa_ext_image_data_get_info)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_data_get_info_fn =
|
||||
image_api.hsa_ext_image_data_get_info_fn =
|
||||
(decltype(::hsa_ext_image_data_get_info)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_create_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_create_fn ==
|
||||
assert(image_api.hsa_ext_image_create_fn ==
|
||||
(decltype(::hsa_ext_image_create)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_create_fn = (decltype(::hsa_ext_image_create)*)ptr;
|
||||
image_api.hsa_ext_image_create_fn = (decltype(::hsa_ext_image_create)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_import_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_import_fn ==
|
||||
assert(image_api.hsa_ext_image_import_fn ==
|
||||
(decltype(::hsa_ext_image_import)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_import_fn = (decltype(::hsa_ext_image_import)*)ptr;
|
||||
image_api.hsa_ext_image_import_fn = (decltype(::hsa_ext_image_import)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_export_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_export_fn ==
|
||||
assert(image_api.hsa_ext_image_export_fn ==
|
||||
(decltype(::hsa_ext_image_export)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_export_fn = (decltype(::hsa_ext_image_export)*)ptr;
|
||||
image_api.hsa_ext_image_export_fn = (decltype(::hsa_ext_image_export)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_copy_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_copy_fn ==
|
||||
assert(image_api.hsa_ext_image_copy_fn ==
|
||||
(decltype(::hsa_ext_image_copy)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_copy_fn = (decltype(::hsa_ext_image_copy)*)ptr;
|
||||
image_api.hsa_ext_image_copy_fn = (decltype(::hsa_ext_image_copy)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_clear_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_clear_fn ==
|
||||
assert(image_api.hsa_ext_image_clear_fn ==
|
||||
(decltype(::hsa_ext_image_clear)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_clear_fn = (decltype(::hsa_ext_image_clear)*)ptr;
|
||||
image_api.hsa_ext_image_clear_fn = (decltype(::hsa_ext_image_clear)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_image_destroy_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_image_destroy_fn ==
|
||||
assert(image_api.hsa_ext_image_destroy_fn ==
|
||||
(decltype(::hsa_ext_image_destroy)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_image_destroy_fn = (decltype(::hsa_ext_image_destroy)*)ptr;
|
||||
image_api.hsa_ext_image_destroy_fn = (decltype(::hsa_ext_image_destroy)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_sampler_create_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_sampler_create_fn ==
|
||||
assert(image_api.hsa_ext_sampler_create_fn ==
|
||||
(decltype(::hsa_ext_sampler_create)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_sampler_create_fn = (decltype(::hsa_ext_sampler_create)*)ptr;
|
||||
image_api.hsa_ext_sampler_create_fn = (decltype(::hsa_ext_sampler_create)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_sampler_destroy_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_ext_sampler_destroy_fn ==
|
||||
assert(image_api.hsa_ext_sampler_destroy_fn ==
|
||||
(decltype(::hsa_ext_sampler_destroy)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_ext_sampler_destroy_fn =
|
||||
image_api.hsa_ext_sampler_destroy_fn =
|
||||
(decltype(::hsa_ext_sampler_destroy)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_amd_image_get_info_max_dim_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_amd_image_get_info_max_dim_fn ==
|
||||
assert(image_api.hsa_amd_image_get_info_max_dim_fn ==
|
||||
(decltype(::hsa_amd_image_get_info_max_dim)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_amd_image_get_info_max_dim_fn =
|
||||
image_api.hsa_amd_image_get_info_max_dim_fn =
|
||||
(decltype(::hsa_amd_image_get_info_max_dim)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_amd_image_create_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(table.hsa_amd_image_create_fn ==
|
||||
(decltype(::hsa_amd_image_create)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
table.hsa_amd_image_create_fn =
|
||||
(decltype(::hsa_amd_image_create)*)ptr;
|
||||
UpdateAmdExtTable(ptr);
|
||||
}
|
||||
|
||||
// Initialize Version of Api Table
|
||||
image_api.version.major_id = HSA_IMAGE_API_TABLE_MAJOR_VERSION;
|
||||
image_api.version.minor_id = sizeof(ImageExtTable);
|
||||
image_api.version.step_id = HSA_IMAGE_API_TABLE_STEP_VERSION;
|
||||
|
||||
// Update private copy of Api table with handle for Image extensions
|
||||
hsa_internal_api_table_.CloneExts(&image_api,
|
||||
core::HsaApiTable::HSA_EXT_IMAGE_API_TABLE_ID);
|
||||
|
||||
core::hsa_internal_api_table_.extension_backup=table;
|
||||
core::hsa_internal_api_table_.table.std_exts_=&core::hsa_internal_api_table_.extension_backup;
|
||||
|
||||
ptr = os::GetExportAddress(lib, "Load");
|
||||
if (ptr != NULL) {
|
||||
((Load_t)ptr)(&core::hsa_internal_api_table_.table);
|
||||
((Load_t)ptr)(&core::hsa_internal_api_table_.hsa_api);
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool ExtensionEntryPoints::LoadFinalizer(std::string library_name) {
|
||||
os::LibHandle lib = os::LoadLib(library_name);
|
||||
if (lib == NULL) {
|
||||
return false;
|
||||
}
|
||||
libs_.push_back(lib);
|
||||
|
||||
void* ptr;
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_create_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(finalizer_api.hsa_ext_program_create_fn ==
|
||||
(decltype(::hsa_ext_program_create)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
finalizer_api.hsa_ext_program_create_fn = (decltype(::hsa_ext_program_create)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_destroy_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(finalizer_api.hsa_ext_program_destroy_fn ==
|
||||
(decltype(::hsa_ext_program_destroy)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
finalizer_api.hsa_ext_program_destroy_fn =
|
||||
(decltype(::hsa_ext_program_destroy)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_add_module_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(finalizer_api.hsa_ext_program_add_module_fn ==
|
||||
(decltype(::hsa_ext_program_add_module)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
finalizer_api.hsa_ext_program_add_module_fn =
|
||||
(decltype(::hsa_ext_program_add_module)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_iterate_modules_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(finalizer_api.hsa_ext_program_iterate_modules_fn ==
|
||||
(decltype(::hsa_ext_program_iterate_modules)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
finalizer_api.hsa_ext_program_iterate_modules_fn =
|
||||
(decltype(::hsa_ext_program_iterate_modules)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_get_info_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(finalizer_api.hsa_ext_program_get_info_fn ==
|
||||
(decltype(::hsa_ext_program_get_info)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
finalizer_api.hsa_ext_program_get_info_fn =
|
||||
(decltype(::hsa_ext_program_get_info)*)ptr;
|
||||
}
|
||||
|
||||
ptr = os::GetExportAddress(lib, "hsa_ext_program_finalize_impl");
|
||||
if (ptr != NULL) {
|
||||
assert(finalizer_api.hsa_ext_program_finalize_fn ==
|
||||
(decltype(::hsa_ext_program_finalize)*)hsa_ext_null &&
|
||||
"Duplicate load of extension import.");
|
||||
finalizer_api.hsa_ext_program_finalize_fn =
|
||||
(decltype(::hsa_ext_program_finalize)*)ptr;
|
||||
}
|
||||
|
||||
// Initialize Version of Api Table
|
||||
finalizer_api.version.major_id = HSA_FINALIZER_API_TABLE_MAJOR_VERSION;
|
||||
finalizer_api.version.minor_id = sizeof(::FinalizerExtTable);
|
||||
finalizer_api.version.step_id = HSA_FINALIZER_API_TABLE_STEP_VERSION;
|
||||
|
||||
// Update handle of table of HSA extensions
|
||||
hsa_internal_api_table_.CloneExts(&finalizer_api,
|
||||
core::HsaApiTable::HSA_EXT_FINALIZER_API_TABLE_ID);
|
||||
|
||||
ptr = os::GetExportAddress(lib, "Load");
|
||||
if (ptr != NULL) {
|
||||
((Load_t)ptr)(&core::hsa_internal_api_table_.hsa_api);
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
||||
} // namespace core
|
||||
|
||||
//---------------------------------------------------------------------------//
|
||||
@@ -386,19 +468,19 @@ hsa_status_t hsa_ext_program_create(
|
||||
hsa_machine_model_t machine_model, hsa_profile_t profile,
|
||||
hsa_default_float_rounding_mode_t default_float_rounding_mode,
|
||||
const char* options, hsa_ext_program_t* program) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.finalizer_api
|
||||
.hsa_ext_program_create_fn(machine_model, profile,
|
||||
default_float_rounding_mode, options, program);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_ext_program_destroy(hsa_ext_program_t program) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.finalizer_api
|
||||
.hsa_ext_program_destroy_fn(program);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_ext_program_add_module(hsa_ext_program_t program,
|
||||
hsa_ext_module_t module) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.finalizer_api
|
||||
.hsa_ext_program_add_module_fn(program, module);
|
||||
}
|
||||
|
||||
@@ -407,14 +489,14 @@ hsa_status_t hsa_ext_program_iterate_modules(
|
||||
hsa_status_t (*callback)(hsa_ext_program_t program, hsa_ext_module_t module,
|
||||
void* data),
|
||||
void* data) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.finalizer_api
|
||||
.hsa_ext_program_iterate_modules_fn(program, callback, data);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_ext_program_get_info(hsa_ext_program_t program,
|
||||
hsa_ext_program_info_t attribute,
|
||||
void* value) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.finalizer_api
|
||||
.hsa_ext_program_get_info_fn(program, attribute, value);
|
||||
}
|
||||
|
||||
@@ -422,7 +504,7 @@ hsa_status_t hsa_ext_program_finalize(
|
||||
hsa_ext_program_t program, hsa_isa_t isa, int32_t call_convention,
|
||||
hsa_ext_control_directives_t control_directives, const char* options,
|
||||
hsa_code_object_type_t code_object_type, hsa_code_object_t* code_object) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.finalizer_api
|
||||
.hsa_ext_program_finalize_fn(program, isa, call_convention,
|
||||
control_directives, options,
|
||||
code_object_type, code_object);
|
||||
@@ -431,7 +513,7 @@ hsa_status_t hsa_ext_program_finalize(
|
||||
hsa_status_t hsa_ext_image_get_capability(
|
||||
hsa_agent_t agent, hsa_ext_image_geometry_t geometry,
|
||||
const hsa_ext_image_format_t* image_format, uint32_t* capability_mask) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_get_capability_fn(agent, geometry, image_format,
|
||||
capability_mask);
|
||||
}
|
||||
@@ -440,7 +522,7 @@ hsa_status_t hsa_ext_image_data_get_info(
|
||||
hsa_agent_t agent, const hsa_ext_image_descriptor_t* image_descriptor,
|
||||
hsa_access_permission_t access_permission,
|
||||
hsa_ext_image_data_info_t* image_data_info) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_data_get_info_fn(agent, image_descriptor,
|
||||
access_permission, image_data_info);
|
||||
}
|
||||
@@ -449,7 +531,7 @@ hsa_status_t hsa_ext_image_create(
|
||||
hsa_agent_t agent, const hsa_ext_image_descriptor_t* image_descriptor,
|
||||
const void* image_data, hsa_access_permission_t access_permission,
|
||||
hsa_ext_image_t* image) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_create_fn(agent, image_descriptor, image_data,
|
||||
access_permission, image);
|
||||
}
|
||||
@@ -458,7 +540,7 @@ hsa_status_t hsa_ext_image_import(hsa_agent_t agent, const void* src_memory,
|
||||
size_t src_row_pitch, size_t src_slice_pitch,
|
||||
hsa_ext_image_t dst_image,
|
||||
const hsa_ext_image_region_t* image_region) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_import_fn(agent, src_memory, src_row_pitch,
|
||||
src_slice_pitch, dst_image, image_region);
|
||||
}
|
||||
@@ -467,7 +549,7 @@ hsa_status_t hsa_ext_image_export(hsa_agent_t agent, hsa_ext_image_t src_image,
|
||||
void* dst_memory, size_t dst_row_pitch,
|
||||
size_t dst_slice_pitch,
|
||||
const hsa_ext_image_region_t* image_region) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_export_fn(agent, src_image, dst_memory, dst_row_pitch,
|
||||
dst_slice_pitch, image_region);
|
||||
}
|
||||
@@ -477,7 +559,7 @@ hsa_status_t hsa_ext_image_copy(hsa_agent_t agent, hsa_ext_image_t src_image,
|
||||
hsa_ext_image_t dst_image,
|
||||
const hsa_dim3_t* dst_offset,
|
||||
const hsa_dim3_t* range) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_copy_fn(agent, src_image, src_offset, dst_image,
|
||||
dst_offset, range);
|
||||
}
|
||||
@@ -485,25 +567,25 @@ hsa_status_t hsa_ext_image_copy(hsa_agent_t agent, hsa_ext_image_t src_image,
|
||||
hsa_status_t hsa_ext_image_clear(hsa_agent_t agent, hsa_ext_image_t image,
|
||||
const void* data,
|
||||
const hsa_ext_image_region_t* image_region) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_clear_fn(agent, image, data, image_region);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_ext_image_destroy(hsa_agent_t agent, hsa_ext_image_t image) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_image_destroy_fn(agent, image);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_ext_sampler_create(
|
||||
hsa_agent_t agent, const hsa_ext_sampler_descriptor_t* sampler_descriptor,
|
||||
hsa_ext_sampler_t* sampler) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_sampler_create_fn(agent, sampler_descriptor, sampler);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_ext_sampler_destroy(hsa_agent_t agent,
|
||||
hsa_ext_sampler_t sampler) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_ext_sampler_destroy_fn(agent, sampler);
|
||||
}
|
||||
|
||||
@@ -511,20 +593,12 @@ hsa_status_t hsa_ext_sampler_destroy(hsa_agent_t agent,
|
||||
// Stubs for internal extension functions
|
||||
//---------------------------------------------------------------------------//
|
||||
|
||||
// Use the function pointer from local instance Image Extension
|
||||
hsa_status_t hsa_amd_image_get_info_max_dim(hsa_agent_t component,
|
||||
hsa_agent_info_t attribute,
|
||||
void* value) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
return core::Runtime::runtime_singleton_->extensions_.image_api
|
||||
.hsa_amd_image_get_info_max_dim_fn(component, attribute, value);
|
||||
}
|
||||
|
||||
hsa_status_t hsa_amd_image_create(
|
||||
hsa_agent_t agent,
|
||||
const hsa_ext_image_descriptor_t *image_descriptor,
|
||||
const hsa_amd_image_descriptor_t *image_layout,
|
||||
const void *image_data,
|
||||
hsa_access_permission_t access_permission,
|
||||
hsa_ext_image_t *image) {
|
||||
return core::Runtime::runtime_singleton_->extensions_.table
|
||||
.hsa_amd_image_create_fn(agent, image_descriptor, image_layout, image_data, access_permission, image);
|
||||
}
|
||||
|
||||
|
||||
@@ -58,6 +58,7 @@
|
||||
#include "core/inc/amd_topology.h"
|
||||
#include "core/inc/signal.h"
|
||||
#include "core/inc/interrupt_signal.h"
|
||||
#include "core/inc/hsa_ext_amd_impl.h"
|
||||
|
||||
#include "core/inc/hsa_api_trace_int.h"
|
||||
|
||||
@@ -528,11 +529,11 @@ hsa_status_t Runtime::GetSystemInfo(hsa_system_info_t attribute, void* value) {
|
||||
case HSA_SYSTEM_INFO_EXTENSIONS:
|
||||
memset(value, 0, sizeof(uint8_t) * 128);
|
||||
|
||||
if (extensions_.table.hsa_ext_program_finalize_fn != NULL) {
|
||||
if (hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
|
||||
*((uint8_t*)value) = 1 << HSA_EXTENSION_FINALIZER;
|
||||
}
|
||||
|
||||
if (extensions_.table.hsa_ext_image_create_fn != NULL) {
|
||||
if (hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
|
||||
*((uint8_t*)value) |= 1 << HSA_EXTENSION_IMAGES;
|
||||
}
|
||||
|
||||
@@ -652,7 +653,7 @@ void Runtime::AsyncEventsLoop(void*) {
|
||||
while (!async_events_control_.exit) {
|
||||
// Wait for a signal
|
||||
hsa_signal_value_t value;
|
||||
uint32_t index = hsa_amd_signal_wait_any(
|
||||
uint32_t index = AMD::hsa_amd_signal_wait_any(
|
||||
uint32_t(async_events_.Size()), &async_events_.signal_[0],
|
||||
&async_events_.cond_[0], &async_events_.value_[0], uint64_t(-1),
|
||||
HSA_WAIT_STATE_BLOCKED, &value);
|
||||
@@ -863,8 +864,12 @@ void Runtime::LoadExtensions() {
|
||||
static const std::string kImageLib[] = {"hsa-ext-image.dll",
|
||||
"libhsa-ext-image.so.1"};
|
||||
#endif
|
||||
extensions_.Load(kFinalizerLib[os_index(os::current_os)]);
|
||||
extensions_.Load(kImageLib[os_index(os::current_os)]);
|
||||
|
||||
// Update Hsa Api Table with handle of Image extension Apis
|
||||
extensions_.LoadFinalizer(kFinalizerLib[os_index(os::current_os)]);
|
||||
|
||||
// Update Hsa Api Table with handle of Finalizer extension Apis
|
||||
extensions_.LoadImage(kImageLib[os_index(os::current_os)]);
|
||||
}
|
||||
|
||||
void Runtime::UnloadExtensions() { extensions_.Unload(); }
|
||||
@@ -920,13 +925,16 @@ static std::vector<std::string> parse_tool_names(std::string tool_names) {
|
||||
}
|
||||
|
||||
void Runtime::LoadTools() {
|
||||
typedef bool (*tool_init_t)(::ApiTable*, uint64_t, uint64_t,
|
||||
typedef bool (*tool_init_t)(::HsaApiTable*, uint64_t, uint64_t,
|
||||
const char* const*);
|
||||
typedef Agent* (*tool_wrap_t)(Agent*);
|
||||
typedef void (*tool_add_t)(Runtime*);
|
||||
|
||||
// Link extensions to API interception
|
||||
hsa_api_table_.LinkExts(&extensions_.table);
|
||||
// Link HSA Extensions for Finalizer and Images for Api interception
|
||||
hsa_api_table_.LinkExts(&extensions_.finalizer_api,
|
||||
core::HsaApiTable::HSA_EXT_FINALIZER_API_TABLE_ID);
|
||||
hsa_api_table_.LinkExts(&extensions_.image_api,
|
||||
core::HsaApiTable::HSA_EXT_IMAGE_API_TABLE_ID);
|
||||
|
||||
// Load tool libs
|
||||
std::string tool_names = flag_.tools_lib_names();
|
||||
@@ -942,7 +950,9 @@ void Runtime::LoadTools() {
|
||||
tool_init_t ld;
|
||||
ld = (tool_init_t)os::GetExportAddress(tool, "OnLoad");
|
||||
if (ld) {
|
||||
if (!ld(&hsa_api_table_.table, 0, failed.size(), &failed[0])) {
|
||||
if (!ld(&hsa_api_table_.hsa_api,
|
||||
hsa_api_table_.hsa_api.version.major_id,
|
||||
failed.size(), &failed[0])) {
|
||||
failed.push_back(names[i].c_str());
|
||||
os::CloseLib(tool);
|
||||
continue;
|
||||
|
||||
@@ -54,13 +54,51 @@
|
||||
#include "inc/hsa_ext_finalize.h"
|
||||
#endif
|
||||
|
||||
struct ExtTable {
|
||||
#include <string.h>
|
||||
#include <assert.h>
|
||||
#include <stddef.h>
|
||||
|
||||
// Major Ids of the Api tables exported by Hsa Core Runtime
|
||||
#define HSA_API_TABLE_MAJOR_VERSION 0x01
|
||||
#define HSA_CORE_API_TABLE_MAJOR_VERSION 0x01
|
||||
#define HSA_AMD_EXT_API_TABLE_MAJOR_VERSION 0x01
|
||||
#define HSA_FINALIZER_API_TABLE_MAJOR_VERSION 0x01
|
||||
#define HSA_IMAGE_API_TABLE_MAJOR_VERSION 0x01
|
||||
|
||||
// Step Ids of the Api tables exported by Hsa Core Runtime
|
||||
#define HSA_API_TABLE_STEP_VERSION 0x00
|
||||
#define HSA_CORE_API_TABLE_STEP_VERSION 0x00
|
||||
#define HSA_AMD_EXT_API_TABLE_STEP_VERSION 0x00
|
||||
#define HSA_FINALIZER_API_TABLE_STEP_VERSION 0x00
|
||||
#define HSA_IMAGE_API_TABLE_STEP_VERSION 0x00
|
||||
|
||||
// Min function used to copy Api Tables
|
||||
static inline uint32_t Min(const uint32_t a, const uint32_t b) {
|
||||
return (a > b) ? b : a;
|
||||
}
|
||||
|
||||
// Structure of Version used to identify an instance of Api table
|
||||
struct ApiTableVersion {
|
||||
uint32_t major_id;
|
||||
uint32_t minor_id;
|
||||
uint32_t step_id;
|
||||
uint32_t reserved;
|
||||
};
|
||||
|
||||
// Table to export HSA Finalizer Extension Apis
|
||||
struct FinalizerExtTable {
|
||||
ApiTableVersion version;
|
||||
decltype(hsa_ext_program_create)* hsa_ext_program_create_fn;
|
||||
decltype(hsa_ext_program_destroy)* hsa_ext_program_destroy_fn;
|
||||
decltype(hsa_ext_program_add_module)* hsa_ext_program_add_module_fn;
|
||||
decltype(hsa_ext_program_iterate_modules)* hsa_ext_program_iterate_modules_fn;
|
||||
decltype(hsa_ext_program_get_info)* hsa_ext_program_get_info_fn;
|
||||
decltype(hsa_ext_program_finalize)* hsa_ext_program_finalize_fn;
|
||||
};
|
||||
|
||||
// Table to export HSA Image Extension Apis
|
||||
struct ImageExtTable {
|
||||
ApiTableVersion version;
|
||||
decltype(hsa_ext_image_get_capability)* hsa_ext_image_get_capability_fn;
|
||||
decltype(hsa_ext_image_data_get_info)* hsa_ext_image_data_get_info_fn;
|
||||
decltype(hsa_ext_image_create)* hsa_ext_image_create_fn;
|
||||
@@ -73,7 +111,40 @@ struct ExtTable {
|
||||
decltype(hsa_ext_sampler_destroy)* hsa_ext_sampler_destroy_fn;
|
||||
};
|
||||
|
||||
struct ApiTable {
|
||||
// Table to export AMD Extension Apis
|
||||
struct AmdExtTable {
|
||||
ApiTableVersion version;
|
||||
decltype(hsa_amd_coherency_get_type)* hsa_amd_coherency_get_type_fn;
|
||||
decltype(hsa_amd_coherency_set_type)* hsa_amd_coherency_set_type_fn;
|
||||
decltype(hsa_amd_profiling_set_profiler_enabled)* hsa_amd_profiling_set_profiler_enabled_fn;
|
||||
decltype(hsa_amd_profiling_async_copy_enable) *hsa_amd_profiling_async_copy_enable_fn;
|
||||
decltype(hsa_amd_profiling_get_dispatch_time)* hsa_amd_profiling_get_dispatch_time_fn;
|
||||
decltype(hsa_amd_profiling_get_async_copy_time) *hsa_amd_profiling_get_async_copy_time_fn;
|
||||
decltype(hsa_amd_profiling_convert_tick_to_system_domain)* hsa_amd_profiling_convert_tick_to_system_domain_fn;
|
||||
decltype(hsa_amd_signal_async_handler)* hsa_amd_signal_async_handler_fn;
|
||||
decltype(hsa_amd_async_function)* hsa_amd_async_function_fn;
|
||||
decltype(hsa_amd_signal_wait_any)* hsa_amd_signal_wait_any_fn;
|
||||
decltype(hsa_amd_queue_cu_set_mask)* hsa_amd_queue_cu_set_mask_fn;
|
||||
decltype(hsa_amd_memory_pool_get_info)* hsa_amd_memory_pool_get_info_fn;
|
||||
decltype(hsa_amd_agent_iterate_memory_pools)* hsa_amd_agent_iterate_memory_pools_fn;
|
||||
decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn;
|
||||
decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn;
|
||||
decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
|
||||
decltype(hsa_amd_agent_memory_pool_get_info)* hsa_amd_agent_memory_pool_get_info_fn;
|
||||
decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn;
|
||||
decltype(hsa_amd_memory_pool_can_migrate)* hsa_amd_memory_pool_can_migrate_fn;
|
||||
decltype(hsa_amd_memory_migrate)* hsa_amd_memory_migrate_fn;
|
||||
decltype(hsa_amd_memory_lock)* hsa_amd_memory_lock_fn;
|
||||
decltype(hsa_amd_memory_unlock)* hsa_amd_memory_unlock_fn;
|
||||
decltype(hsa_amd_memory_fill)* hsa_amd_memory_fill_fn;
|
||||
decltype(hsa_amd_interop_map_buffer)* hsa_amd_interop_map_buffer_fn;
|
||||
decltype(hsa_amd_interop_unmap_buffer)* hsa_amd_interop_unmap_buffer_fn;
|
||||
decltype(::hsa_amd_image_create)* hsa_amd_image_create_fn;
|
||||
};
|
||||
|
||||
// Table to export HSA Core Runtime Apis
|
||||
struct CoreApiTable {
|
||||
ApiTableVersion version;
|
||||
decltype(hsa_init)* hsa_init_fn;
|
||||
decltype(hsa_shut_down)* hsa_shut_down_fn;
|
||||
decltype(hsa_system_get_info)* hsa_system_get_info_fn;
|
||||
@@ -170,8 +241,126 @@ struct ApiTable {
|
||||
decltype(hsa_executable_symbol_get_info)* hsa_executable_symbol_get_info_fn;
|
||||
decltype(hsa_executable_iterate_symbols)* hsa_executable_iterate_symbols_fn;
|
||||
decltype(hsa_status_string)* hsa_status_string_fn;
|
||||
|
||||
ExtTable* std_exts_;
|
||||
};
|
||||
|
||||
// Table to export HSA Apis from Core Runtime, Amd Extensions
|
||||
// Finalizer and Images
|
||||
struct HsaApiTable {
|
||||
|
||||
// Version of Hsa Api Table
|
||||
ApiTableVersion version;
|
||||
|
||||
// Table of function pointers to HSA Core Runtime
|
||||
CoreApiTable* core_;
|
||||
|
||||
// Table of function pointers to AMD extensions
|
||||
AmdExtTable* amd_ext_;
|
||||
|
||||
// Table of function pointers to HSA Finalizer Extension
|
||||
FinalizerExtTable* finalizer_ext_;
|
||||
|
||||
// Table of function pointers to HSA Image Extension
|
||||
ImageExtTable* image_ext_;
|
||||
|
||||
};
|
||||
|
||||
// Structure containing instances of different api tables
|
||||
struct HsaApiTableContainer {
|
||||
HsaApiTable root;
|
||||
CoreApiTable core;
|
||||
AmdExtTable amd_ext;
|
||||
FinalizerExtTable finalizer_ext;
|
||||
ImageExtTable image_ext;
|
||||
|
||||
// Default initialization of a container instance
|
||||
HsaApiTableContainer() {
|
||||
root.version.major_id = HSA_API_TABLE_MAJOR_VERSION;
|
||||
root.version.minor_id = sizeof(HsaApiTable);
|
||||
root.version.step_id = HSA_API_TABLE_STEP_VERSION;
|
||||
|
||||
core.version.major_id = HSA_CORE_API_TABLE_MAJOR_VERSION;
|
||||
core.version.minor_id = sizeof(CoreApiTable);
|
||||
core.version.step_id = HSA_CORE_API_TABLE_STEP_VERSION;
|
||||
root.core_ = &core;
|
||||
|
||||
amd_ext.version.major_id = HSA_AMD_EXT_API_TABLE_MAJOR_VERSION;
|
||||
amd_ext.version.minor_id = sizeof(AmdExtTable);
|
||||
amd_ext.version.step_id = HSA_AMD_EXT_API_TABLE_STEP_VERSION;
|
||||
root.amd_ext_ = &amd_ext;
|
||||
|
||||
finalizer_ext.version.major_id = HSA_FINALIZER_API_TABLE_MAJOR_VERSION;
|
||||
finalizer_ext.version.minor_id = sizeof(FinalizerExtTable);
|
||||
finalizer_ext.version.step_id = HSA_FINALIZER_API_TABLE_STEP_VERSION;
|
||||
root.finalizer_ext_ = & finalizer_ext;
|
||||
|
||||
image_ext.version.major_id = HSA_IMAGE_API_TABLE_MAJOR_VERSION;
|
||||
image_ext.version.minor_id = sizeof(ImageExtTable);
|
||||
image_ext.version.step_id = HSA_IMAGE_API_TABLE_STEP_VERSION;
|
||||
root.image_ext_ = &image_ext;
|
||||
}
|
||||
};
|
||||
|
||||
// Api to copy function pointers of a table
|
||||
static
|
||||
void inline copyApi(void* src, void* dest, size_t size) {
|
||||
memcpy((char*)src + sizeof(ApiTableVersion),
|
||||
(char*)dest + sizeof(ApiTableVersion),
|
||||
(size - sizeof(ApiTableVersion)));
|
||||
}
|
||||
|
||||
// Copy constructor for all Api tables. The function assumes the
|
||||
// user has initialized an instance of tables container correctly
|
||||
// for the Major, Minor and Stepping Ids of Root and Child Api tables.
|
||||
// The function will overwrite the value of Minor Id by taking the
|
||||
// minimum of source and destination parameters. It will also overwrite
|
||||
// the stepping Id with value from source parameter.
|
||||
static const
|
||||
void inline copyTables(const HsaApiTable* src, HsaApiTableContainer* dest) {
|
||||
|
||||
// Verify Major Id of source and destination tables are valid
|
||||
assert(dest->root.version.major_id == src->version.major_id);
|
||||
assert(dest->core.version.major_id == src->core_->version.major_id);
|
||||
assert(dest->amd_ext.version.major_id == src->amd_ext_->version.major_id);
|
||||
assert(dest->finalizer_ext.version.major_id == src->finalizer_ext_->version.major_id);
|
||||
assert(dest->image_ext.version.major_id == src->image_ext_->version.major_id);
|
||||
|
||||
// Initialize the stepping id and minor id of root table. For the
|
||||
// minor id which encodes struct size, take the minimum of source
|
||||
// and destination parameters
|
||||
dest->root.version.step_id = src->version.step_id;
|
||||
dest->root.version.minor_id = Min(dest->root.version.minor_id, src->version.minor_id);
|
||||
|
||||
// Copy the Core Api table
|
||||
size_t size = dest->root.version.minor_id;
|
||||
if (size > offsetof(HsaApiTable, core_)) {
|
||||
dest->core.version.step_id = src->core_->version.step_id;
|
||||
dest->core.version.minor_id = Min(dest->core.version.minor_id,
|
||||
src->core_->version.minor_id);
|
||||
copyApi(&dest->core, src->core_, dest->core.version.minor_id);
|
||||
}
|
||||
|
||||
// Copy the Amd Ext Api table
|
||||
if (size > offsetof(HsaApiTable, amd_ext_)) {
|
||||
dest->amd_ext.version.step_id = src->amd_ext_->version.step_id;
|
||||
dest->amd_ext.version.minor_id = Min(dest->core.version.minor_id,
|
||||
src->amd_ext_->version.minor_id);
|
||||
copyApi(&dest->amd_ext, src->amd_ext_, dest->amd_ext.version.minor_id);
|
||||
}
|
||||
|
||||
// Copy the Finalizer Ext Api table
|
||||
if (size > offsetof(HsaApiTable, finalizer_ext_)) {
|
||||
dest->finalizer_ext.version.step_id = src->finalizer_ext_->version.step_id;
|
||||
dest->finalizer_ext.version.minor_id = Min(dest->core.version.minor_id,
|
||||
src->finalizer_ext_->version.minor_id);
|
||||
copyApi(&dest->finalizer_ext, src->finalizer_ext_, dest->finalizer_ext.version.minor_id);
|
||||
}
|
||||
|
||||
// Copy the Image Ext Api table
|
||||
if (size > offsetof(HsaApiTable, image_ext_)) {
|
||||
dest->image_ext.version.step_id = src->image_ext_->version.step_id;
|
||||
dest->image_ext.version.minor_id = Min(dest->core.version.minor_id,
|
||||
src->image_ext_->version.minor_id);
|
||||
copyApi(&dest->image_ext, src->image_ext_, dest->image_ext.version.minor_id);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
In neuem Issue referenzieren
Einen Benutzer sperren