From d260c22467dd8a1a18d853d2a6786649ec0f13f4 Mon Sep 17 00:00:00 2001 From: Ramesh Errabolu Date: Fri, 10 Jun 2016 17:31:13 -0500 Subject: [PATCH] Export Amd Extension APIs including support for Version Control Change-Id: I8c03cbd4049e8115ae00d51f193b9c31ac941f21 [ROCm/ROCR-Runtime commit: 95dc97da7b96a31104ddde8aaa1abdd2c219e6ec] --- .../core/common/hsa_table_interface.cpp | 423 ++++++++++++++---- .../hsa-runtime/core/inc/hsa_api_trace_int.h | 29 +- .../hsa-runtime/core/inc/hsa_ext_amd_impl.h | 186 ++++++++ .../hsa-runtime/core/inc/hsa_ext_interface.h | 28 +- .../core/inc/hsa_table_interface.h | 4 +- .../core/runtime/amd_aql_queue.cpp | 3 +- .../core/runtime/amd_gpu_agent.cpp | 6 +- .../runtime/hsa-runtime/core/runtime/hsa.cpp | 107 ++--- .../core/runtime/hsa_api_trace.cpp | 324 +++++++++----- .../hsa-runtime/core/runtime/hsa_ext_amd.cpp | 50 ++- .../core/runtime/hsa_ext_interface.cpp | 342 ++++++++------ .../hsa-runtime/core/runtime/runtime.cpp | 28 +- .../runtime/hsa-runtime/inc/hsa_api_trace.h | 197 +++++++- 13 files changed, 1278 insertions(+), 449 deletions(-) create mode 100755 projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/common/hsa_table_interface.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/common/hsa_table_interface.cpp index ffbe749a98..13154820f6 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/common/hsa_table_interface.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/common/hsa_table_interface.cpp @@ -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); +} + diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h index cc9a638a9d..769dbed2e0 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h @@ -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 diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h new file mode 100755 index 0000000000..54f8e34588 --- /dev/null +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h @@ -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 diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_interface.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_interface.h index 3645c23fbb..236a165c73 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_interface.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_ext_interface.h @@ -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 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); }; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_table_interface.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_table_interface.h index 236ef41c7d..99a1280d8f 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_table_interface.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_table_interface.h @@ -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(); diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 7821a9da69..81b08e8e1e 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -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; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index 32ca97cacb..d96fd414d9 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -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; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp index 3f69ec74c6..f5b49d1ad2 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -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(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(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(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(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(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(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(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(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; } //---------------------------------------------------------------------------// diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp index ca0b401925..9fe3823a5b 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp @@ -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 + 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; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp index 45a64a8c5f..a31b5a6c35 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp @@ -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(&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 diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp index 3aa9f5c04e..c8d8bf541d 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp @@ -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 +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); -} + diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp index 9abca46d08..0d397b2240 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -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 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; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_api_trace.h b/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_api_trace.h index ee7e63b9ec..40d443de1d 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_api_trace.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_api_trace.h @@ -54,13 +54,51 @@ #include "inc/hsa_ext_finalize.h" #endif -struct ExtTable { +#include +#include +#include + +// 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