From 7e2179da7ba864c9819a2def6e6bede995124ce4 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Tue, 23 Aug 2016 05:30:22 -0500 Subject: [PATCH] Update clang-format file to clang-format v3.8. Format HSA v1.1 core updates. Change-Id: I540b5c0e5b3ec7522b09c2e070167812b3f17769 [ROCm/ROCR-Runtime commit: 54f1311e01bee8951f33b5ac219b5db38da6ddd6] --- .../core/common/hsa_table_interface.cpp | 207 ++++----- .../runtime/hsa-runtime/core/inc/agent.h | 3 +- .../hsa-runtime/core/inc/amd_cpu_agent.h | 3 +- .../hsa-runtime/core/inc/amd_gpu_agent.h | 3 +- .../runtime/hsa-runtime/core/inc/cache.h | 66 ++- .../hsa-runtime/core/inc/hsa_internal.h | 159 ++++--- .../runtime/hsa-runtime/core/inc/queue.h | 10 +- .../runtime/hsa-runtime/core/inc/signal.h | 23 +- .../core/runtime/amd_blit_kernel.cpp | 21 +- .../core/runtime/amd_cpu_agent.cpp | 19 +- .../core/runtime/amd_gpu_agent.cpp | 53 +-- .../hsa-runtime/core/runtime/cache.cpp | 22 +- .../runtime/hsa-runtime/core/runtime/hsa.cpp | 405 ++++++++---------- .../core/runtime/hsa_api_trace.cpp | 34 +- .../hsa-runtime/core/runtime/runtime.cpp | 34 +- .../hsa-runtime/core/runtime/signal.cpp | 24 +- .../runtime/hsa-runtime/inc/hsa_api_trace.h | 2 +- 17 files changed, 507 insertions(+), 581 deletions(-) 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 a0370c672b..8df54eebfa 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 @@ -67,8 +67,7 @@ hsa_status_t HSA_API return coreApiTable->hsa_system_get_info_fn(attribute, value); } -hsa_status_t HSA_API hsa_extension_get_name(uint16_t extension, const char** name) -{ +hsa_status_t HSA_API hsa_extension_get_name(uint16_t extension, const char** name) { return coreApiTable->hsa_extension_get_name_fn(extension, name); } @@ -79,23 +78,24 @@ hsa_status_t HSA_API extension, version_major, version_minor, result); } -hsa_status_t HSA_API - hsa_system_major_extension_supported(uint16_t extension, uint16_t version_major, - uint16_t *version_minor, bool* result) { - return coreApiTable->hsa_system_major_extension_supported_fn( - extension, version_major, version_minor, result); +hsa_status_t HSA_API hsa_system_major_extension_supported(uint16_t extension, + uint16_t version_major, + uint16_t* version_minor, bool* result) { + return coreApiTable->hsa_system_major_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) { +hsa_status_t HSA_API hsa_system_get_extension_table(uint16_t extension, uint16_t version_major, + uint16_t version_minor, void* table) { return coreApiTable->hsa_system_get_extension_table_fn( extension, version_major, version_minor, table); } -hsa_status_t HSA_API hsa_system_get_major_extension_table(uint16_t extension, uint16_t version_major, size_t table_length, void *table) -{ - return coreApiTable->hsa_system_get_major_extension_table_fn(extension, version_major, table_length, table); +hsa_status_t HSA_API hsa_system_get_major_extension_table(uint16_t extension, + uint16_t version_major, + size_t table_length, void* table) { + return coreApiTable->hsa_system_get_major_extension_table_fn(extension, version_major, + table_length, table); } hsa_status_t HSA_API @@ -116,13 +116,13 @@ hsa_status_t HSA_API hsa_agent_get_exception_policies(hsa_agent_t agent, return coreApiTable->hsa_agent_get_exception_policies_fn(agent, profile, mask); } -hsa_status_t HSA_API hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, void *value) -{ +hsa_status_t HSA_API hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, + void* value) { return coreApiTable->hsa_cache_get_info_fn(cache, attribute, value); } -hsa_status_t HSA_API hsa_agent_iterate_caches(hsa_agent_t agent, hsa_status_t (*callback )(hsa_cache_t cache, void *data), void *value) -{ +hsa_status_t HSA_API hsa_agent_iterate_caches( + hsa_agent_t agent, hsa_status_t (*callback)(hsa_cache_t cache, void* data), void* value) { return coreApiTable->hsa_agent_iterate_caches_fn(agent, callback, value); } @@ -134,9 +134,11 @@ hsa_status_t HSA_API extension, agent, version_major, version_minor, result); } -hsa_status_t HSA_API hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t agent, uint16_t version_major, uint16_t *version_minor, bool *result) -{ - return coreApiTable->hsa_agent_major_extension_supported_fn(extension, agent, version_major, version_minor, result); +hsa_status_t HSA_API hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t agent, + uint16_t version_major, + uint16_t* version_minor, bool* result) { + return coreApiTable->hsa_agent_major_extension_supported_fn(extension, agent, version_major, + version_minor, result); } hsa_status_t HSA_API @@ -187,23 +189,18 @@ void HSA_API hsa_queue_store_write_index_relaxed(const hsa_queue_t* queue, return coreApiTable->hsa_queue_store_write_index_relaxed_fn(queue, value); } -void HSA_API hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, - uint64_t value) { +void HSA_API hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { return coreApiTable->hsa_queue_store_write_index_screlease_fn(queue, value); } -uint64_t HSA_API hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, - uint64_t expected, - uint64_t value) { - return coreApiTable->hsa_queue_cas_write_index_scacq_screl_fn(queue, expected, - value); +uint64_t HSA_API hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, uint64_t expected, + uint64_t value) { + return coreApiTable->hsa_queue_cas_write_index_scacq_screl_fn(queue, expected, value); } -uint64_t HSA_API hsa_queue_cas_write_index_scacquire(const hsa_queue_t* queue, - uint64_t expected, - uint64_t value) { - return coreApiTable->hsa_queue_cas_write_index_scacquire_fn(queue, expected, - value); +uint64_t HSA_API hsa_queue_cas_write_index_scacquire(const hsa_queue_t* queue, uint64_t expected, + uint64_t value) { + return coreApiTable->hsa_queue_cas_write_index_scacquire_fn(queue, expected, value); } uint64_t HSA_API hsa_queue_cas_write_index_relaxed(const hsa_queue_t* queue, @@ -213,20 +210,16 @@ uint64_t HSA_API hsa_queue_cas_write_index_relaxed(const hsa_queue_t* queue, value); } -uint64_t HSA_API hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, - uint64_t expected, - uint64_t value) { - return coreApiTable->hsa_queue_cas_write_index_screlease_fn(queue, expected, - value); +uint64_t HSA_API hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, uint64_t expected, + uint64_t value) { + return coreApiTable->hsa_queue_cas_write_index_screlease_fn(queue, expected, value); } -uint64_t HSA_API hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, - uint64_t value) { +uint64_t HSA_API hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, uint64_t value) { return coreApiTable->hsa_queue_add_write_index_scacq_screl_fn(queue, value); } -uint64_t HSA_API hsa_queue_add_write_index_scacquire(const hsa_queue_t* queue, - uint64_t value) { +uint64_t HSA_API hsa_queue_add_write_index_scacquire(const hsa_queue_t* queue, uint64_t value) { return coreApiTable->hsa_queue_add_write_index_scacquire_fn(queue, value); } @@ -235,8 +228,7 @@ uint64_t HSA_API hsa_queue_add_write_index_relaxed(const hsa_queue_t* queue, return coreApiTable->hsa_queue_add_write_index_relaxed_fn(queue, value); } -uint64_t HSA_API hsa_queue_add_write_index_screlease(const hsa_queue_t* queue, - uint64_t value) { +uint64_t HSA_API hsa_queue_add_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { return coreApiTable->hsa_queue_add_write_index_screlease_fn(queue, value); } @@ -245,8 +237,7 @@ void HSA_API hsa_queue_store_read_index_relaxed(const hsa_queue_t* queue, return coreApiTable->hsa_queue_store_read_index_relaxed_fn(queue, value); } -void HSA_API hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, - uint64_t value) { +void HSA_API hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, uint64_t value) { return coreApiTable->hsa_queue_store_read_index_screlease_fn(queue, value); } @@ -312,18 +303,15 @@ void HSA_API return coreApiTable->hsa_signal_store_relaxed_fn(signal, value); } -void HSA_API - hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_store_screlease_fn(signal, value); } -void HSA_API hsa_signal_silent_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value) -{ +void HSA_API hsa_signal_silent_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_silent_store_relaxed_fn(signal, value); } -void HSA_API hsa_signal_silent_store_screlease(hsa_signal_t signal, hsa_signal_value_t value) -{ +void HSA_API hsa_signal_silent_store_screlease(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_silent_store_screlease_fn(signal, value); } @@ -337,34 +325,44 @@ hsa_signal_value_t HSA_API signal, condition, compare_value, timeout_hint, wait_expectancy_hint); } -hsa_signal_value_t HSA_API - hsa_signal_wait_scacquire(hsa_signal_t signal, - hsa_signal_condition_t condition, - hsa_signal_value_t compare_value, - uint64_t timeout_hint, - hsa_wait_state_t wait_expectancy_hint) { - return coreApiTable->hsa_signal_wait_scacquire_fn( - signal, condition, compare_value, timeout_hint, wait_expectancy_hint); +hsa_signal_value_t HSA_API hsa_signal_wait_scacquire(hsa_signal_t signal, + hsa_signal_condition_t condition, + hsa_signal_value_t compare_value, + uint64_t timeout_hint, + hsa_wait_state_t wait_expectancy_hint) { + return coreApiTable->hsa_signal_wait_scacquire_fn(signal, condition, compare_value, timeout_hint, + wait_expectancy_hint); } -hsa_status_t HSA_API hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t *signals, uint32_t num_consumers, const hsa_agent_t *consumers, hsa_signal_group_t *signal_group) -{ - return coreApiTable->hsa_signal_group_create_fn(num_signals, signals, num_consumers, consumers, signal_group); +hsa_status_t HSA_API hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t* signals, + uint32_t num_consumers, const hsa_agent_t* consumers, + hsa_signal_group_t* signal_group) { + return coreApiTable->hsa_signal_group_create_fn(num_signals, signals, num_consumers, consumers, + signal_group); } -hsa_status_t HSA_API hsa_signal_group_destroy(hsa_signal_group_t signal_group) -{ +hsa_status_t HSA_API hsa_signal_group_destroy(hsa_signal_group_t signal_group) { return coreApiTable->hsa_signal_group_destroy_fn(signal_group); } -hsa_status_t HSA_API hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, const hsa_signal_condition_t *conditions, const hsa_signal_value_t *compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t *signal, hsa_signal_value_t *value) -{ - return coreApiTable->hsa_signal_group_wait_any_relaxed_fn(signal_group, conditions, compare_values, wait_state_hint, signal, value); +hsa_status_t HSA_API hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, + const hsa_signal_condition_t* conditions, + const hsa_signal_value_t* compare_values, + hsa_wait_state_t wait_state_hint, + hsa_signal_t* signal, + hsa_signal_value_t* value) { + return coreApiTable->hsa_signal_group_wait_any_relaxed_fn( + signal_group, conditions, compare_values, wait_state_hint, signal, value); } -hsa_status_t HSA_API hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, const hsa_signal_condition_t *conditions, const hsa_signal_value_t *compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t *signal, hsa_signal_value_t *value) -{ - return coreApiTable->hsa_signal_group_wait_any_scacquire_fn(signal_group, conditions, compare_values, wait_state_hint, signal, value); +hsa_status_t HSA_API hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, + const hsa_signal_condition_t* conditions, + const hsa_signal_value_t* compare_values, + hsa_wait_state_t wait_state_hint, + hsa_signal_t* signal, + hsa_signal_value_t* value) { + return coreApiTable->hsa_signal_group_wait_any_scacquire_fn( + signal_group, conditions, compare_values, wait_state_hint, signal, value); } void HSA_API @@ -372,18 +370,15 @@ void HSA_API return coreApiTable->hsa_signal_and_relaxed_fn(signal, value); } -void HSA_API - hsa_signal_and_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_and_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_and_scacquire_fn(signal, value); } -void HSA_API - hsa_signal_and_screlease(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_and_screlease(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_and_screlease_fn(signal, value); } -void HSA_API - hsa_signal_and_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_and_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_and_scacq_screl_fn(signal, value); } @@ -392,18 +387,15 @@ void HSA_API return coreApiTable->hsa_signal_or_relaxed_fn(signal, value); } -void HSA_API - hsa_signal_or_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_or_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_or_scacquire_fn(signal, value); } -void HSA_API - hsa_signal_or_screlease(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_or_screlease(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_or_screlease_fn(signal, value); } -void HSA_API - hsa_signal_or_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_or_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_or_scacq_screl_fn(signal, value); } @@ -412,18 +404,15 @@ void HSA_API return coreApiTable->hsa_signal_xor_relaxed_fn(signal, value); } -void HSA_API - hsa_signal_xor_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_xor_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_xor_scacquire_fn(signal, value); } -void HSA_API - hsa_signal_xor_screlease(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_xor_screlease(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_xor_screlease_fn(signal, value); } -void HSA_API - hsa_signal_xor_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_xor_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_xor_scacq_screl_fn(signal, value); } @@ -432,18 +421,15 @@ void HSA_API return coreApiTable->hsa_signal_add_relaxed_fn(signal, value); } -void HSA_API - hsa_signal_add_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_add_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_add_scacquire_fn(signal, value); } -void HSA_API - hsa_signal_add_screlease(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_add_screlease(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_add_screlease_fn(signal, value); } -void HSA_API - hsa_signal_add_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_add_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_add_scacq_screl_fn(signal, value); } @@ -452,18 +438,15 @@ void HSA_API return coreApiTable->hsa_signal_subtract_relaxed_fn(signal, value); } -void HSA_API - hsa_signal_subtract_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_subtract_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_subtract_scacquire_fn(signal, value); } -void HSA_API - hsa_signal_subtract_screlease(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_subtract_screlease(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_subtract_screlease_fn(signal, value); } -void HSA_API - hsa_signal_subtract_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { +void HSA_API hsa_signal_subtract_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { return coreApiTable->hsa_signal_subtract_scacq_screl_fn(signal, value); } @@ -472,18 +455,18 @@ hsa_signal_value_t HSA_API return coreApiTable->hsa_signal_exchange_relaxed_fn(signal, value); } -hsa_signal_value_t HSA_API - hsa_signal_exchange_scacquire(hsa_signal_t signal, hsa_signal_value_t value) { +hsa_signal_value_t HSA_API hsa_signal_exchange_scacquire(hsa_signal_t signal, + hsa_signal_value_t value) { return coreApiTable->hsa_signal_exchange_scacquire_fn(signal, value); } -hsa_signal_value_t HSA_API - hsa_signal_exchange_screlease(hsa_signal_t signal, hsa_signal_value_t value) { +hsa_signal_value_t HSA_API hsa_signal_exchange_screlease(hsa_signal_t signal, + hsa_signal_value_t value) { return coreApiTable->hsa_signal_exchange_screlease_fn(signal, value); } -hsa_signal_value_t HSA_API - hsa_signal_exchange_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value) { +hsa_signal_value_t HSA_API hsa_signal_exchange_scacq_screl(hsa_signal_t signal, + hsa_signal_value_t value) { return coreApiTable->hsa_signal_exchange_scacq_screl_fn(signal, value); } @@ -494,20 +477,20 @@ hsa_signal_value_t HSA_API hsa_signal_cas_relaxed(hsa_signal_t signal, } hsa_signal_value_t HSA_API hsa_signal_cas_scacquire(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value) { + hsa_signal_value_t expected, + hsa_signal_value_t value) { return coreApiTable->hsa_signal_cas_scacquire_fn(signal, expected, value); } hsa_signal_value_t HSA_API hsa_signal_cas_screlease(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value) { + hsa_signal_value_t expected, + hsa_signal_value_t value) { return coreApiTable->hsa_signal_cas_screlease_fn(signal, expected, value); } hsa_signal_value_t HSA_API hsa_signal_cas_scacq_screl(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value) { + hsa_signal_value_t expected, + hsa_signal_value_t value) { return coreApiTable->hsa_signal_cas_scacq_screl_fn(signal, expected, value); } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/agent.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/agent.h index c0d41b827a..b1c55f6627 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/agent.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/agent.h @@ -191,7 +191,8 @@ class Agent : public Checked<0xF6BC25EB17E6F917> { void* data) const = 0; // @brief Invoke the callback for each cache useable by this agent. - virtual hsa_status_t IterateCache(hsa_status_t (*callback )(hsa_cache_t cache, void *data), void* data) const = 0; + virtual hsa_status_t IterateCache(hsa_status_t (*callback)(hsa_cache_t cache, void* data), + void* data) const = 0; // @brief Create queue. // diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_cpu_agent.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_cpu_agent.h index 19762bb87d..e7c1a03855 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_cpu_agent.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_cpu_agent.h @@ -90,7 +90,8 @@ class CpuAgent : public core::Agent { void* data) const override; // @brief Override from core::Agent. - hsa_status_t IterateCache(hsa_status_t (*callback )(hsa_cache_t cache, void *data), void* value) const override; + hsa_status_t IterateCache(hsa_status_t (*callback)(hsa_cache_t cache, void* data), + void* value) const override; // @brief Override from core::Agent. hsa_status_t GetInfo(hsa_agent_info_t attribute, void* value) const override; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index a7c4f08570..3cee37d44f 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -224,7 +224,8 @@ class GpuAgent : public GpuAgentInt { void* data) const override; // @brief Override from core::Agent. - hsa_status_t IterateCache(hsa_status_t (*callback )(hsa_cache_t cache, void *data), void* value) const override; + hsa_status_t IterateCache(hsa_status_t (*callback)(hsa_cache_t cache, void* data), + void* value) const override; // @brief Override from core::Agent. hsa_status_t DmaCopy(void* dst, const void* src, size_t size) override; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/cache.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/cache.h index 1f0b6479be..9e7f7d12a3 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/cache.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/cache.h @@ -2,24 +2,24 @@ // // 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 @@ -29,7 +29,7 @@ // 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 @@ -49,38 +49,34 @@ #include #include -namespace core -{ +namespace core { - class Cache : public Checked<0x39A6C7AD3F135B06> - { - public: - static __forceinline hsa_cache_t Convert(const Cache* cache) - { - const hsa_cache_t handle={ static_cast(reinterpret_cast(cache)) }; - return handle; - } - static __forceinline Cache* Convert(const hsa_cache_t cache) - { - return reinterpret_cast(static_cast(cache.handle)); - } +class Cache : public Checked<0x39A6C7AD3F135B06> { + public: + static __forceinline hsa_cache_t Convert(const Cache* cache) { + const hsa_cache_t handle = {static_cast(reinterpret_cast(cache))}; + return handle; + } + static __forceinline Cache* Convert(const hsa_cache_t cache) { + return reinterpret_cast(static_cast(cache.handle)); + } - Cache(const std::string& name, uint8_t level, uint32_t size) : name_(name), level_(level), size_(size) {} + Cache(const std::string& name, uint8_t level, uint32_t size) + : name_(name), level_(level), size_(size) {} - Cache(std::string&& name, uint8_t level, uint32_t size) : name_(std::move(name)), level_(level), size_(size) {} + Cache(std::string&& name, uint8_t level, uint32_t size) + : name_(std::move(name)), level_(level), size_(size) {} - hsa_status_t GetInfo(hsa_cache_info_t attribute, void* value); + hsa_status_t GetInfo(hsa_cache_info_t attribute, void* value); - private: - - std::string name_; - uint32_t level_; - uint32_t size_; - - // Forbid copying and moving of this object - DISALLOW_COPY_AND_ASSIGN(Cache); - }; + private: + std::string name_; + uint32_t level_; + uint32_t size_; + // Forbid copying and moving of this object + DISALLOW_COPY_AND_ASSIGN(Cache); +}; } -#endif // HSA_RUNTIME_CORE_INC_CACHE_H +#endif // HSA_RUNTIME_CORE_INC_CACHE_H diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_internal.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_internal.h index 887a0ea22e..b6c9f26c4a 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_internal.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/hsa_internal.h @@ -53,16 +53,18 @@ namespace HSA hsa_status_t HSA_API hsa_shut_down(); hsa_status_t HSA_API hsa_system_get_info(hsa_system_info_t attribute, void *value); - hsa_status_t HSA_API hsa_extension_get_name(uint16_t extension, const char **name); + hsa_status_t HSA_API hsa_extension_get_name(uint16_t extension, const char** name); hsa_status_t HSA_API hsa_system_extension_supported(uint16_t extension, uint16_t version_major, - uint16_t version_minor, bool* result); - hsa_status_t HSA_API - hsa_system_major_extension_supported(uint16_t extension, uint16_t version_major, - uint16_t *version_minor, bool *result); + uint16_t version_minor, bool* result); + hsa_status_t HSA_API hsa_system_major_extension_supported(uint16_t extension, + uint16_t version_major, + uint16_t* version_minor, bool* result); hsa_status_t HSA_API hsa_system_get_extension_table(uint16_t extension, uint16_t version_major, uint16_t version_minor, void *table); - hsa_status_t HSA_API hsa_system_get_major_extension_table(uint16_t extension, uint16_t version_major, size_t table_length, void *table); + hsa_status_t HSA_API hsa_system_get_major_extension_table(uint16_t extension, + uint16_t version_major, + size_t table_length, void* table); hsa_status_t HSA_API hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, void *data), void *data); @@ -72,13 +74,17 @@ namespace HSA hsa_status_t HSA_API hsa_agent_get_exception_policies(hsa_agent_t agent, hsa_profile_t profile, uint16_t *mask); - hsa_status_t HSA_API hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, void *value); - hsa_status_t HSA_API hsa_agent_iterate_caches(hsa_agent_t agent, hsa_status_t (*callback )(hsa_cache_t cache, void *data), void *value); + hsa_status_t HSA_API hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, + void* value); + hsa_status_t HSA_API hsa_agent_iterate_caches( + hsa_agent_t agent, hsa_status_t (*callback)(hsa_cache_t cache, void* data), void* value); 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); - hsa_status_t HSA_API hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t agent, uint16_t version_major, uint16_t *version_minor, bool *result); + hsa_status_t HSA_API hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t agent, + uint16_t version_major, + uint16_t* version_minor, bool* result); hsa_status_t HSA_API hsa_queue_create(hsa_agent_t agent, uint32_t size, hsa_queue_type_t type, void (*callback)(hsa_status_t status, hsa_queue_t *source, @@ -91,38 +97,30 @@ namespace HSA hsa_signal_t completion_signal, hsa_queue_t **queue); hsa_status_t HSA_API hsa_queue_destroy(hsa_queue_t *queue); hsa_status_t HSA_API hsa_queue_inactivate(hsa_queue_t *queue); - uint64_t HSA_API hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue); + uint64_t HSA_API hsa_queue_load_read_index_scacquire(const hsa_queue_t* queue); uint64_t HSA_API hsa_queue_load_read_index_relaxed(const hsa_queue_t *queue); - uint64_t HSA_API hsa_queue_load_write_index_scacquire(const hsa_queue_t *queue); + uint64_t HSA_API hsa_queue_load_write_index_scacquire(const hsa_queue_t* queue); uint64_t HSA_API hsa_queue_load_write_index_relaxed(const hsa_queue_t *queue); void HSA_API hsa_queue_store_write_index_relaxed(const hsa_queue_t *queue, uint64_t value); - void HSA_API hsa_queue_store_write_index_screlease(const hsa_queue_t *queue, - uint64_t value); - uint64_t HSA_API hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t *queue, - uint64_t expected, - uint64_t value); - uint64_t HSA_API hsa_queue_cas_write_index_scacquire(const hsa_queue_t *queue, - uint64_t expected, - uint64_t value); + void HSA_API hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, uint64_t value); + uint64_t HSA_API hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, + uint64_t expected, uint64_t value); + uint64_t HSA_API hsa_queue_cas_write_index_scacquire(const hsa_queue_t* queue, uint64_t expected, + uint64_t value); uint64_t HSA_API hsa_queue_cas_write_index_relaxed(const hsa_queue_t *queue, uint64_t expected, uint64_t value); - uint64_t HSA_API hsa_queue_cas_write_index_screlease(const hsa_queue_t *queue, - uint64_t expected, - uint64_t value); - uint64_t HSA_API - hsa_queue_add_write_index_scacq_screl(const hsa_queue_t *queue, uint64_t value); - uint64_t HSA_API - hsa_queue_add_write_index_scacquire(const hsa_queue_t *queue, uint64_t value); + uint64_t HSA_API hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, uint64_t expected, + uint64_t value); + uint64_t HSA_API hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, uint64_t value); + uint64_t HSA_API hsa_queue_add_write_index_scacquire(const hsa_queue_t* queue, uint64_t value); uint64_t HSA_API hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue, uint64_t value); - uint64_t HSA_API - hsa_queue_add_write_index_screlease(const hsa_queue_t *queue, uint64_t value); + uint64_t HSA_API hsa_queue_add_write_index_screlease(const hsa_queue_t* queue, uint64_t value); void HSA_API hsa_queue_store_read_index_relaxed(const hsa_queue_t *queue, uint64_t value); - void HSA_API hsa_queue_store_read_index_screlease(const hsa_queue_t *queue, - uint64_t value); + void HSA_API hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, uint64_t 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); @@ -145,8 +143,7 @@ namespace HSA hsa_signal_value_t HSA_API hsa_signal_load_scacquire(hsa_signal_t signal); void HSA_API hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value); void HSA_API hsa_signal_silent_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value); void HSA_API hsa_signal_silent_store_screlease(hsa_signal_t signal, hsa_signal_value_t value); hsa_signal_value_t HSA_API @@ -155,76 +152,72 @@ namespace HSA hsa_signal_value_t compare_value, uint64_t timeout_hint, hsa_wait_state_t wait_expectancy_hint); - hsa_signal_value_t HSA_API - hsa_signal_wait_scacquire(hsa_signal_t signal, - hsa_signal_condition_t condition, - hsa_signal_value_t compare_value, - uint64_t timeout_hint, - hsa_wait_state_t wait_expectancy_hint); - hsa_status_t HSA_API hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t *signals, uint32_t num_consumers, const hsa_agent_t *consumers, hsa_signal_group_t *signal_group); + hsa_signal_value_t HSA_API hsa_signal_wait_scacquire(hsa_signal_t signal, + hsa_signal_condition_t condition, + hsa_signal_value_t compare_value, + uint64_t timeout_hint, + hsa_wait_state_t wait_expectancy_hint); + hsa_status_t HSA_API hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t* signals, + uint32_t num_consumers, const hsa_agent_t* consumers, + hsa_signal_group_t* signal_group); hsa_status_t HSA_API hsa_signal_group_destroy(hsa_signal_group_t signal_group); - hsa_status_t HSA_API hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, const hsa_signal_condition_t *conditions, const hsa_signal_value_t *compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t *signal, hsa_signal_value_t *value); - hsa_status_t HSA_API hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, const hsa_signal_condition_t *conditions, const hsa_signal_value_t *compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t *signal, hsa_signal_value_t *value); + hsa_status_t HSA_API hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, + const hsa_signal_condition_t* conditions, + const hsa_signal_value_t* compare_values, + hsa_wait_state_t wait_state_hint, + hsa_signal_t* signal, + hsa_signal_value_t* value); + hsa_status_t HSA_API hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, + const hsa_signal_condition_t* conditions, + const hsa_signal_value_t* compare_values, + hsa_wait_state_t wait_state_hint, + hsa_signal_t* signal, + hsa_signal_value_t* value); void HSA_API hsa_signal_and_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_and_scacquire(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_and_screlease(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_and_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_and_scacquire(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_and_screlease(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_and_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); void HSA_API hsa_signal_or_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_or_scacquire(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_or_screlease(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_or_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_or_scacquire(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_or_screlease(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_or_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); void HSA_API hsa_signal_xor_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_xor_scacquire(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_xor_screlease(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_xor_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_xor_scacquire(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_xor_screlease(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_xor_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); void HSA_API hsa_signal_add_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_add_scacquire(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_add_screlease(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_add_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_add_scacquire(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_add_screlease(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_add_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); void HSA_API hsa_signal_subtract_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_subtract_scacquire(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_subtract_screlease(hsa_signal_t signal, hsa_signal_value_t value); - void HSA_API - hsa_signal_subtract_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_subtract_scacquire(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_subtract_screlease(hsa_signal_t signal, hsa_signal_value_t value); + void HSA_API hsa_signal_subtract_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); hsa_signal_value_t HSA_API hsa_signal_exchange_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - hsa_signal_value_t HSA_API - hsa_signal_exchange_scacquire(hsa_signal_t signal, hsa_signal_value_t value); - hsa_signal_value_t HSA_API - hsa_signal_exchange_screlease(hsa_signal_t signal, hsa_signal_value_t value); - hsa_signal_value_t HSA_API - hsa_signal_exchange_scacq_screl(hsa_signal_t signal, hsa_signal_value_t value); + hsa_signal_value_t HSA_API hsa_signal_exchange_scacquire(hsa_signal_t signal, + hsa_signal_value_t value); + hsa_signal_value_t HSA_API hsa_signal_exchange_screlease(hsa_signal_t signal, + hsa_signal_value_t value); + hsa_signal_value_t HSA_API hsa_signal_exchange_scacq_screl(hsa_signal_t signal, + hsa_signal_value_t 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); hsa_signal_value_t HSA_API hsa_signal_cas_scacquire(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value); + hsa_signal_value_t expected, + hsa_signal_value_t value); hsa_signal_value_t HSA_API hsa_signal_cas_screlease(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value); + hsa_signal_value_t expected, + hsa_signal_value_t value); hsa_signal_value_t HSA_API hsa_signal_cas_scacq_screl(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value); + hsa_signal_value_t expected, + hsa_signal_value_t value); //===--- Instruction Set Architecture -----------------------------------===// diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h index 0c3aceca1a..402843d7c5 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h @@ -89,12 +89,10 @@ struct AqlPacket { string << "type: " << type_names[type] << "\nbarrier: " << ((dispatch.header >> HSA_PACKET_HEADER_BARRIER) & ((1 << HSA_PACKET_HEADER_WIDTH_BARRIER) - 1)) - << "\nacquire: " - << ((dispatch.header >> HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) & - ((1 << HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE) - 1)) - << "\nrelease: " - << ((dispatch.header >> HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE) & - ((1 << HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE) - 1)); + << "\nacquire: " << ((dispatch.header >> HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) & + ((1 << HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE) - 1)) + << "\nrelease: " << ((dispatch.header >> HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE) & + ((1 << HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE) - 1)); if (type == HSA_PACKET_TYPE_KERNEL_DISPATCH) { string << "\nDim: " << dispatch.setup diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/signal.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/signal.h index 171fd8afce..21e61bdcae 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/signal.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/signal.h @@ -210,9 +210,8 @@ class Signal : public Checked<0x71FCCA6A3D5D5276>, /// Returns the index of a satisfied signal. Returns -1 on timeout and /// errors. static uint32_t WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals, - const hsa_signal_condition_t* conds, - const hsa_signal_value_t* values, uint64_t timeout_hint, - hsa_wait_state_t wait_hint, + const hsa_signal_condition_t* conds, const hsa_signal_value_t* values, + uint64_t timeout_hint, hsa_wait_state_t wait_hint, hsa_signal_value_t* satisfying_value); __forceinline bool IsType(rtti_t id) { return _IsA(id); } @@ -277,24 +276,28 @@ static_assert( sizeof(hsa_signal_handle[2]) == sizeof(hsa_signal_t[2]), "hsa_signal_handle and hsa_signal_t must have identical binary layout."); -class SignalGroup : public Checked<0xBD35DDDD578F091> -{ -public: - static __forceinline hsa_signal_group_t Convert(SignalGroup* group) { const hsa_signal_group_t handle = { static_cast(reinterpret_cast(group)) }; return handle; } - static __forceinline SignalGroup* Convert(hsa_signal_group_t group) { return reinterpret_cast(static_cast(group.handle)); } +class SignalGroup : public Checked<0xBD35DDDD578F091> { + public: + static __forceinline hsa_signal_group_t Convert(SignalGroup* group) { + const hsa_signal_group_t handle = {static_cast(reinterpret_cast(group))}; + return handle; + } + static __forceinline SignalGroup* Convert(hsa_signal_group_t group) { + return reinterpret_cast(static_cast(group.handle)); + } SignalGroup(uint32_t num_signals, const hsa_signal_t* signals); ~SignalGroup() { delete[] signals; } bool IsValid() const { - if (CheckedType::IsValid() && signals!=NULL) return true; + if (CheckedType::IsValid() && signals != NULL) return true; return false; } const hsa_signal_t* List() const { return signals; } uint32_t Count() const { return count; } -private: + private: hsa_signal_t* signals; const uint32_t count; DISALLOW_COPY_AND_ASSIGN(SignalGroup); diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp index e8f0f9f070..679fce594e 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp @@ -603,9 +603,8 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(void* dst, const void* src, } // Wait for the packet to finish. - if (HSA::hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_LT, - 1, uint64_t(-1), - HSA_WAIT_STATE_ACTIVE) != 0) { + if (HSA::hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1), + HSA_WAIT_STATE_ACTIVE) != 0) { // Signal wait returned unexpected value. return HSA_STATUS_ERROR; } @@ -624,8 +623,7 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand( uint64_t write_index_temp = write_index; // Insert barrier packets to handle dependent signals. - const uint16_t kBarrierPacketHeader = - (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) | + const uint16_t kBarrierPacketHeader = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) | (HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); @@ -766,9 +764,8 @@ hsa_status_t BlitKernel::SubmitLinearFillCommand(void* ptr, uint32_t value, ReleaseWriteIndex(write_index, 1); // Wait for the packet to finish. - if (HSA::hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_LT, - 1, uint64_t(-1), - HSA_WAIT_STATE_ACTIVE) != 0) { + if (HSA::hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1), + HSA_WAIT_STATE_ACTIVE) != 0) { // Signal wait returned unexpected value. return HSA_STATUS_ERROR; } @@ -810,8 +807,7 @@ hsa_status_t BlitKernel::FenceRelease(uint64_t write_index, hsa_fence_scope_t fence) { // This function is not thread safe. - const uint16_t kBarrierPacketHeader = - (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) | + const uint16_t kBarrierPacketHeader = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) | (fence << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); @@ -842,9 +838,8 @@ hsa_status_t BlitKernel::FenceRelease(uint64_t write_index, ReleaseWriteIndex(write_index, num_copy_packet + 1); // Wait for the packet to finish. - if (HSA::hsa_signal_wait_scacquire(packet.completion_signal, - HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1), - HSA_WAIT_STATE_ACTIVE) != 0) { + if (HSA::hsa_signal_wait_scacquire(packet.completion_signal, HSA_SIGNAL_CONDITION_LT, 1, + uint64_t(-1), HSA_WAIT_STATE_ACTIVE) != 0) { // Signal wait returned unexpected value. return HSA_STATUS_ERROR; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp index 5bfc0378ac..c84e272a82 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp @@ -126,14 +126,15 @@ void CpuAgent::InitCacheList() { } } - //Update cache objects + // Update cache objects caches_.clear(); caches_.resize(cache_props_.size()); char name[64]; GetInfo(HSA_AGENT_INFO_NAME, name); - std::string deviceName=name; - for(size_t i=0; i struct ValidityError; -template <> -struct ValidityError { +template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_SIGNAL }; }; -template <> -struct ValidityError { +template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP }; }; -template <> -struct ValidityError { +template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_AGENT }; }; -template <> -struct ValidityError { +template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_REGION }; }; -template <> -struct ValidityError { +template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_QUEUE }; }; -template <> -struct ValidityError { +template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_CACHE }; }; -template <> -struct ValidityError { +template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_ISA }; }; -template -struct ValidityError { +template struct ValidityError { enum { kValue = ValidityError::kValue }; }; @@ -212,12 +204,10 @@ hsa_status_t return core::Runtime::runtime_singleton_->GetSystemInfo(attribute, value); } -hsa_status_t hsa_extension_get_name(uint16_t extension, const char** name) -{ +hsa_status_t hsa_extension_get_name(uint16_t extension, const char** name) { IS_OPEN(); IS_BAD_PTR(name); - switch(extension) - { + switch (extension) { case HSA_EXTENSION_FINALIZER: *name = "HSA_EXTENSION_FINALIZER"; break; @@ -248,13 +238,16 @@ hsa_status_t uint16_t version_minor, bool* result) { IS_OPEN(); - if ((extension > HSA_EXTENSION_STD_LAST && (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || result == NULL) { + if ((extension > HSA_EXTENSION_STD_LAST && + (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || + result == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } *result = false; - if(extension==HSA_EXTENSION_PERFORMANCE_COUNTERS || extension == HSA_EXTENSION_PROFILING_EVENTS) + if (extension == HSA_EXTENSION_PERFORMANCE_COUNTERS || + extension == HSA_EXTENSION_PROFILING_EVENTS) return HSA_STATUS_SUCCESS; uint16_t system_version_major = 0; @@ -272,108 +265,98 @@ hsa_status_t return HSA_STATUS_SUCCESS; } -hsa_status_t - hsa_system_major_extension_supported(uint16_t extension, uint16_t version_major, - uint16_t *version_minor, bool* result) { +hsa_status_t hsa_system_major_extension_supported(uint16_t extension, uint16_t version_major, + uint16_t* version_minor, bool* result) { IS_OPEN(); IS_BAD_PTR(version_minor); IS_BAD_PTR(result); - if ((extension == HSA_EXTENSION_IMAGES) && (version_major==1)) { - *version_minor=0; - *result=true; + if ((extension == HSA_EXTENSION_IMAGES) && (version_major == 1)) { + *version_minor = 0; + *result = true; return HSA_STATUS_SUCCESS; } - if ((extension == HSA_EXTENSION_FINALIZER) && (version_major==1)) { - *version_minor=0; - *result=true; + if ((extension == HSA_EXTENSION_FINALIZER) && (version_major == 1)) { + *version_minor = 0; + *result = true; return HSA_STATUS_SUCCESS; } - if ((extension == HSA_EXTENSION_AMD_LOADER) && (version_major==1)) { - *version_minor=0; - *result=true; + if ((extension == HSA_EXTENSION_AMD_LOADER) && (version_major == 1)) { + *version_minor = 0; + *result = true; return HSA_STATUS_SUCCESS; } - *result=false; + *result = false; return HSA_STATUS_SUCCESS; } -static size_t get_extension_table_length(uint16_t extension, uint16_t major, uint16_t minor) -{ - //Table to convert from major/minor to major/length - struct sizes_t - { +static size_t get_extension_table_length(uint16_t extension, uint16_t major, uint16_t minor) { + // Table to convert from major/minor to major/length + struct sizes_t { std::string name; size_t size; }; - static sizes_t sizes[] = { - "hsa_ext_images_1_00_pfn_t", sizeof(hsa_ext_images_1_00_pfn_t), - "hsa_ext_finalizer_1_00_pfn_t", sizeof(hsa_ext_finalizer_1_00_pfn_t), - "hsa_ven_amd_loader_1_00_pfn_t", sizeof(hsa_ven_amd_loader_1_00_pfn_t) - }; - static const size_t num_tables = sizeof(sizes)/sizeof(sizes_t); + static sizes_t sizes[] = {"hsa_ext_images_1_00_pfn_t", sizeof(hsa_ext_images_1_00_pfn_t), + "hsa_ext_finalizer_1_00_pfn_t", sizeof(hsa_ext_finalizer_1_00_pfn_t), + "hsa_ven_amd_loader_1_00_pfn_t", sizeof(hsa_ven_amd_loader_1_00_pfn_t)}; + static const size_t num_tables = sizeof(sizes) / sizeof(sizes_t); - if(minor>99) - return 0; + if (minor > 99) return 0; std::string name; - switch(extension) - { - case HSA_EXTENSION_FINALIZER: - name = "hsa_ext_finalizer_"; - break; - case HSA_EXTENSION_IMAGES: - name="hsa_ext_images_"; - break; - //case HSA_EXTENSION_PERFORMANCE_COUNTERS: - // name = "hsa_ext_perf_counter_"; - // break; - //case HSA_EXTENSION_PROFILING_EVENTS: - // name = "hsa_ext_profiling_event_"; - // break; - //case HSA_EXTENSION_AMD_PROFILER: - // name = "hsa_ven_amd_profiler_"; - // break; - case HSA_EXTENSION_AMD_LOADER: - name = "hsa_ven_amd_loader_"; - break; - default: - return 0; + switch (extension) { + case HSA_EXTENSION_FINALIZER: + name = "hsa_ext_finalizer_"; + break; + case HSA_EXTENSION_IMAGES: + name = "hsa_ext_images_"; + break; + // case HSA_EXTENSION_PERFORMANCE_COUNTERS: + // name = "hsa_ext_perf_counter_"; + // break; + // case HSA_EXTENSION_PROFILING_EVENTS: + // name = "hsa_ext_profiling_event_"; + // break; + // case HSA_EXTENSION_AMD_PROFILER: + // name = "hsa_ven_amd_profiler_"; + // break; + case HSA_EXTENSION_AMD_LOADER: + name = "hsa_ven_amd_loader_"; + break; + default: + return 0; } char buff[3]; sprintf(buff, "%02u", minor); - name+=std::to_string(major)+"_"+buff+"_pfn_t"; + name += std::to_string(major) + "_" + buff + "_pfn_t"; - for(size_t i=0; i1) - return HSA_STATUS_ERROR; + if (version_major > 1) return HSA_STATUS_ERROR; hsa_ext_images_1_00_pfn_t ext_table; ext_table.hsa_ext_image_clear = hsa_ext_image_clear; ext_table.hsa_ext_image_copy = hsa_ext_image_copy; @@ -387,13 +370,12 @@ hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t v ext_table.hsa_ext_sampler_destroy = hsa_ext_sampler_destroy; memcpy(table, &ext_table, Min(sizeof(ext_table), table_length)); - + return HSA_STATUS_SUCCESS; } if (extension == HSA_EXTENSION_FINALIZER) { - if(version_major>1) - return HSA_STATUS_ERROR; + if (version_major > 1) return HSA_STATUS_ERROR; hsa_ext_finalizer_1_00_pfn_t ext_table; ext_table.hsa_ext_program_add_module = hsa_ext_program_add_module; ext_table.hsa_ext_program_create = hsa_ext_program_create; @@ -408,13 +390,11 @@ hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t v } if (extension == HSA_EXTENSION_AMD_LOADER) { - if(version_major>1) - return HSA_STATUS_ERROR; + if (version_major > 1) return HSA_STATUS_ERROR; hsa_ven_amd_loader_1_00_pfn_t ext_table; - ext_table.hsa_ven_amd_loader_query_host_address = - hsa_ven_amd_loader_query_host_address; + ext_table.hsa_ven_amd_loader_query_host_address = hsa_ven_amd_loader_query_host_address; ext_table.hsa_ven_amd_loader_query_segment_descriptors = - hsa_ven_amd_loader_query_segment_descriptors; + hsa_ven_amd_loader_query_segment_descriptors; memcpy(table, &ext_table, Min(sizeof(ext_table), table_length)); @@ -458,8 +438,7 @@ hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent_handle, return HSA_STATUS_SUCCESS; } -hsa_status_t hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, void *value) -{ +hsa_status_t hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, void* value) { IS_OPEN(); core::Cache* Cache = core::Cache::Convert(cache); IS_VALID(Cache); @@ -467,10 +446,11 @@ hsa_status_t hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, v return Cache->GetInfo(attribute, value); } -hsa_status_t hsa_agent_iterate_caches(hsa_agent_t agent_handle, hsa_status_t (*callback )(hsa_cache_t cache, void *data), void *data) -{ +hsa_status_t hsa_agent_iterate_caches(hsa_agent_t agent_handle, + hsa_status_t (*callback)(hsa_cache_t cache, void* data), + void* data) { IS_OPEN(); - const core::Agent* agent=core::Agent::Convert(agent_handle); + const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); IS_BAD_PTR(callback); return agent->IterateCache(callback, data); @@ -482,7 +462,9 @@ hsa_status_t uint16_t version_minor, bool* result) { IS_OPEN(); - if ((extension > HSA_EXTENSION_STD_LAST && (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || result == NULL) { + if ((extension > HSA_EXTENSION_STD_LAST && + (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || + result == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } @@ -508,11 +490,14 @@ hsa_status_t return HSA_STATUS_SUCCESS; } -hsa_status_t hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t agent_handle, uint16_t version_major, uint16_t *version_minor, bool *result) -{ +hsa_status_t hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t agent_handle, + uint16_t version_major, uint16_t* version_minor, + bool* result) { IS_OPEN(); - if ((extension > HSA_EXTENSION_STD_LAST && (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || result == NULL) { + if ((extension > HSA_EXTENSION_STD_LAST && + (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || + result == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } @@ -523,8 +508,7 @@ hsa_status_t hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t if (agent->device_type() == core::Agent::kAmdGpuDevice) { uint16_t agent_version_major = 0; - hsa_status_t status = - agent->GetInfo(HSA_AGENT_INFO_VERSION_MAJOR, &agent_version_major); + hsa_status_t status = agent->GetInfo(HSA_AGENT_INFO_VERSION_MAJOR, &agent_version_major); assert(status == HSA_STATUS_SUCCESS); if (version_major <= agent_version_major) { @@ -715,8 +699,7 @@ void hsa_queue_store_read_index_relaxed(const hsa_queue_t* queue, /// @param queue Pointer to the queue whose read index is being updated /// /// @param value Value of new read index -void hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, - uint64_t value) { +void hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreReadIndexRelease(value); @@ -739,8 +722,7 @@ void hsa_queue_store_write_index_relaxed(const hsa_queue_t* queue, /// @param queue Pointer to the queue whose write index is being updated /// /// @param value Value of new write index -void hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, - uint64_t value) { +void hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreWriteIndexRelease(value); @@ -756,9 +738,8 @@ void hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, /// @param value Value of new write index /// /// @return uint64_t Value of write index before the update -uint64_t hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, - uint64_t expected, - uint64_t value) { +uint64_t hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, uint64_t expected, + uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcqRel(expected, value); @@ -774,9 +755,8 @@ uint64_t hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, /// @param value Value of new write index /// /// @return uint64_t Value of write index before the update -uint64_t hsa_queue_cas_write_index_scacquire(const hsa_queue_t* queue, - uint64_t expected, - uint64_t value) { +uint64_t hsa_queue_cas_write_index_scacquire(const hsa_queue_t* queue, uint64_t expected, + uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcquire(expected, value); @@ -810,9 +790,8 @@ uint64_t hsa_queue_cas_write_index_relaxed(const hsa_queue_t* queue, /// @param value Value of new write index /// /// @return uint64_t Value of write index before the update -uint64_t hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, - uint64_t expected, - uint64_t value) { +uint64_t hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, uint64_t expected, + uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexRelease(expected, value); @@ -826,8 +805,7 @@ uint64_t hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, /// @param value Value to add to write index /// /// @return uint64_t Value of write index before the update -uint64_t hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, - uint64_t value) { +uint64_t hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcqRel(value); @@ -840,8 +818,7 @@ uint64_t hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, /// @param value Value to add to write index /// /// @return uint64_t Value of write index before the update -uint64_t hsa_queue_add_write_index_scacquire(const hsa_queue_t* queue, - uint64_t value) { +uint64_t hsa_queue_add_write_index_scacquire(const hsa_queue_t* queue, uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcquire(value); @@ -868,8 +845,7 @@ uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t* queue, /// @param value Value to add to write index /// /// @return uint64_t Value of write index before the update -uint64_t hsa_queue_add_write_index_screlease(const hsa_queue_t* queue, - uint64_t value) { +uint64_t hsa_queue_add_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexRelease(value); @@ -1060,8 +1036,7 @@ void hsa_signal_store_relaxed(hsa_signal_t hsa_signal, signal->StoreRelaxed(value); } -void hsa_signal_store_screlease(hsa_signal_t hsa_signal, - hsa_signal_value_t value) { +void hsa_signal_store_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->StoreRelease(value); @@ -1079,40 +1054,35 @@ hsa_signal_value_t wait_state_hint); } -hsa_signal_value_t - hsa_signal_wait_scacquire(hsa_signal_t hsa_signal, - hsa_signal_condition_t condition, - hsa_signal_value_t compare_value, - uint64_t timeout_hint, - hsa_wait_state_t wait_state_hint) { +hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t hsa_signal, + hsa_signal_condition_t condition, + hsa_signal_value_t compare_value, + uint64_t timeout_hint, + hsa_wait_state_t wait_state_hint) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->WaitAcquire(condition, compare_value, timeout_hint, wait_state_hint); } -hsa_status_t hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t *signals, uint32_t num_consumers, const hsa_agent_t *consumers, hsa_signal_group_t *signal_group) -{ +hsa_status_t hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t* signals, + uint32_t num_consumers, const hsa_agent_t* consumers, + hsa_signal_group_t* signal_group) { IS_OPEN(); - if(num_signals==0) - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - for (uint i = 0; i < num_signals; i++) - IS_VALID(core::Signal::Convert(signals[i])); - for (uint i = 0; i < num_consumers; i++) - IS_VALID(core::Agent::Convert(consumers[i])); - core::SignalGroup* group=new core::SignalGroup(num_signals, signals); + if (num_signals == 0) return HSA_STATUS_ERROR_INVALID_ARGUMENT; + for (uint i = 0; i < num_signals; i++) IS_VALID(core::Signal::Convert(signals[i])); + for (uint i = 0; i < num_consumers; i++) IS_VALID(core::Agent::Convert(consumers[i])); + core::SignalGroup* group = new core::SignalGroup(num_signals, signals); CHECK_ALLOC(group); - if(!group->IsValid()) - { + if (!group->IsValid()) { delete group; return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } - *signal_group=core::SignalGroup::Convert(group); + *signal_group = core::SignalGroup::Convert(group); return HSA_STATUS_SUCCESS; } -hsa_status_t hsa_signal_group_destroy(hsa_signal_group_t signal_group) -{ +hsa_status_t hsa_signal_group_destroy(hsa_signal_group_t signal_group) { IS_OPEN(); core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); @@ -1120,21 +1090,30 @@ hsa_status_t hsa_signal_group_destroy(hsa_signal_group_t signal_group) return HSA_STATUS_SUCCESS; } -hsa_status_t hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, const hsa_signal_condition_t *conditions, const hsa_signal_value_t *compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t *signal, hsa_signal_value_t *value) -{ +hsa_status_t hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, + const hsa_signal_condition_t* conditions, + const hsa_signal_value_t* compare_values, + hsa_wait_state_t wait_state_hint, + hsa_signal_t* signal, hsa_signal_value_t* value) { IS_OPEN(); const core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); - const uint32_t index = AMD::hsa_amd_signal_wait_any(group->Count(), const_cast(group->List()), const_cast(conditions), const_cast(compare_values), uint64_t(-1), wait_state_hint, value); - if(index>=group->Count()) - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - *signal=group->List()[index]; + const uint32_t index = AMD::hsa_amd_signal_wait_any( + group->Count(), const_cast(group->List()), + const_cast(conditions), + const_cast(compare_values), uint64_t(-1), wait_state_hint, value); + if (index >= group->Count()) return HSA_STATUS_ERROR_INVALID_ARGUMENT; + *signal = group->List()[index]; return HSA_STATUS_SUCCESS; } -hsa_status_t hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, const hsa_signal_condition_t *conditions, const hsa_signal_value_t *compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t *signal, hsa_signal_value_t *value) -{ - hsa_status_t ret=HSA::hsa_signal_group_wait_any_relaxed(signal_group, conditions, compare_values, wait_state_hint, signal, value); +hsa_status_t hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, + const hsa_signal_condition_t* conditions, + const hsa_signal_value_t* compare_values, + hsa_wait_state_t wait_state_hint, + hsa_signal_t* signal, hsa_signal_value_t* value) { + hsa_status_t ret = HSA::hsa_signal_group_wait_any_relaxed( + signal_group, conditions, compare_values, wait_state_hint, signal, value); std::atomic_thread_fence(std::memory_order_acquire); return ret; } @@ -1146,22 +1125,19 @@ void signal->AndRelaxed(value); } -void - hsa_signal_and_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_and_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcquire(value); } -void - hsa_signal_and_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_and_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndRelease(value); } -void - hsa_signal_and_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_and_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcqRel(value); @@ -1174,22 +1150,19 @@ void signal->OrRelaxed(value); } -void - hsa_signal_or_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_or_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcquire(value); } -void - hsa_signal_or_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_or_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrRelease(value); } -void - hsa_signal_or_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_or_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcqRel(value); @@ -1202,22 +1175,19 @@ void signal->XorRelaxed(value); } -void - hsa_signal_xor_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_xor_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcquire(value); } -void - hsa_signal_xor_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_xor_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorRelease(value); } -void - hsa_signal_xor_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_xor_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcqRel(value); @@ -1230,22 +1200,19 @@ void return signal->AddRelaxed(value); } -void - hsa_signal_add_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_add_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcquire(value); } -void - hsa_signal_add_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_add_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddRelease(value); } -void - hsa_signal_add_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_add_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcqRel(value); @@ -1258,22 +1225,19 @@ void hsa_signal_subtract_relaxed(hsa_signal_t hsa_signal, signal->SubRelaxed(value); } -void hsa_signal_subtract_scacquire(hsa_signal_t hsa_signal, - hsa_signal_value_t value) { +void hsa_signal_subtract_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcquire(value); } -void hsa_signal_subtract_screlease(hsa_signal_t hsa_signal, - hsa_signal_value_t value) { +void hsa_signal_subtract_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubRelease(value); } -void hsa_signal_subtract_scacq_screl(hsa_signal_t hsa_signal, - hsa_signal_value_t value) { +void hsa_signal_subtract_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcqRel(value); @@ -1287,25 +1251,22 @@ hsa_signal_value_t return signal->ExchRelaxed(value); } -hsa_signal_value_t - hsa_signal_exchange_scacquire(hsa_signal_t hsa_signal, - hsa_signal_value_t value) { +hsa_signal_value_t hsa_signal_exchange_scacquire(hsa_signal_t hsa_signal, + hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcquire(value); } -hsa_signal_value_t - hsa_signal_exchange_screlease(hsa_signal_t hsa_signal, - hsa_signal_value_t value) { +hsa_signal_value_t hsa_signal_exchange_screlease(hsa_signal_t hsa_signal, + hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchRelease(value); } -hsa_signal_value_t - hsa_signal_exchange_scacq_screl(hsa_signal_t hsa_signal, - hsa_signal_value_t value) { +hsa_signal_value_t hsa_signal_exchange_scacq_screl(hsa_signal_t hsa_signal, + hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcqRel(value); @@ -1319,25 +1280,22 @@ hsa_signal_value_t hsa_signal_cas_relaxed(hsa_signal_t hsa_signal, return signal->CasRelaxed(expected, value); } -hsa_signal_value_t hsa_signal_cas_scacquire(hsa_signal_t hsa_signal, - hsa_signal_value_t expected, - hsa_signal_value_t value) { +hsa_signal_value_t hsa_signal_cas_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t expected, + hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcquire(expected, value); } -hsa_signal_value_t hsa_signal_cas_screlease(hsa_signal_t hsa_signal, - hsa_signal_value_t expected, - hsa_signal_value_t value) { +hsa_signal_value_t hsa_signal_cas_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t expected, + hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasRelease(expected, value); } -hsa_signal_value_t hsa_signal_cas_scacq_screl(hsa_signal_t hsa_signal, - hsa_signal_value_t expected, - hsa_signal_value_t value) { +hsa_signal_value_t hsa_signal_cas_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t expected, + hsa_signal_value_t value) { core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcqRel(expected, value); @@ -2335,46 +2293,53 @@ hsa_status_t hsa_status_string( *status_string = "HSA_STATUS_ERROR_FROZEN_EXECUTABLE: The executable is frozen."; break; case HSA_STATUS_ERROR_INVALID_SYMBOL_NAME: - *status_string = "HSA_STATUS_ERROR_INVALID_SYMBOL_NAME: There is no symbol with the given name."; + *status_string = + "HSA_STATUS_ERROR_INVALID_SYMBOL_NAME: There is no symbol with the given name."; break; case HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED: - *status_string = "HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED: The variable is already defined."; + *status_string = + "HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED: The variable is already defined."; break; case HSA_STATUS_ERROR_VARIABLE_UNDEFINED: *status_string = "HSA_STATUS_ERROR_VARIABLE_UNDEFINED: The variable is undefined."; break; - case HSA_STATUS_ERROR_INVALID_CODE_SYMBOL : + case HSA_STATUS_ERROR_INVALID_CODE_SYMBOL: *status_string = "HSA_STATUS_ERROR_INVALID_CODE_SYMBOL: The code object symbol is invalid."; break; - case HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL : - *status_string = "HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL: The executable symbol is invalid."; + case HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL: + *status_string = + "HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL: The executable symbol is invalid."; break; - case HSA_STATUS_ERROR_INVALID_FILE : + case HSA_STATUS_ERROR_INVALID_FILE: *status_string = "HSA_STATUS_ERROR_INVALID_FILE: The file descriptor is invalid."; break; - case HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER : - *status_string = "HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: *The code object reader is invalid."; + case HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: + *status_string = + "HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: *The code object reader is invalid."; break; - case HSA_STATUS_ERROR_INVALID_CACHE : + case HSA_STATUS_ERROR_INVALID_CACHE: *status_string = "HSA_STATUS_ERROR_INVALID_CACHE: *The cache is invalid."; break; - case HSA_STATUS_ERROR_INVALID_WAVEFRONT : + case HSA_STATUS_ERROR_INVALID_WAVEFRONT: *status_string = "HSA_STATUS_ERROR_INVALID_WAVEFRONT: *The wavefront is invalid."; break; - case HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP : + case HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP: *status_string = "HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP: *The signal group is invalid."; break; - case HSA_STATUS_ERROR_INVALID_RUNTIME_STATE : - *status_string = "HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: *The HSA runtime is not in the configuration state."; + case HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: + *status_string = + "HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: *The HSA runtime is not in the configuration " + "state."; break; case HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED: - *status_string = "HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED: Image " - "format is not supported."; + *status_string = + "HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED: Image " + "format is not supported."; break; case HSA_EXT_STATUS_ERROR_IMAGE_SIZE_UNSUPPORTED: *status_string = - "HSA_EXT_STATUS_ERROR_IMAGE_SIZE_UNSUPPORTED: Image size " - "is not supported."; + "HSA_EXT_STATUS_ERROR_IMAGE_SIZE_UNSUPPORTED: Image size " + "is not supported."; break; case HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED: *status_string = "Image pitch is not supported or invalid."; 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 c64194affd..a2d56d0326 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 @@ -146,38 +146,28 @@ void HsaApiTable::UpdateCore() { 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_scacquire_fn = - HSA::hsa_queue_load_read_index_scacquire; + core_api.hsa_queue_load_read_index_scacquire_fn = HSA::hsa_queue_load_read_index_scacquire; core_api.hsa_queue_load_read_index_relaxed_fn = HSA::hsa_queue_load_read_index_relaxed; - core_api.hsa_queue_load_write_index_scacquire_fn = - HSA::hsa_queue_load_write_index_scacquire; + core_api.hsa_queue_load_write_index_scacquire_fn = HSA::hsa_queue_load_write_index_scacquire; core_api.hsa_queue_load_write_index_relaxed_fn = HSA::hsa_queue_load_write_index_relaxed; core_api.hsa_queue_store_write_index_relaxed_fn = HSA::hsa_queue_store_write_index_relaxed; - core_api.hsa_queue_store_write_index_screlease_fn = - HSA::hsa_queue_store_write_index_screlease; - core_api.hsa_queue_cas_write_index_scacq_screl_fn = - HSA::hsa_queue_cas_write_index_scacq_screl; - core_api.hsa_queue_cas_write_index_scacquire_fn = - HSA::hsa_queue_cas_write_index_scacquire; + core_api.hsa_queue_store_write_index_screlease_fn = HSA::hsa_queue_store_write_index_screlease; + core_api.hsa_queue_cas_write_index_scacq_screl_fn = HSA::hsa_queue_cas_write_index_scacq_screl; + core_api.hsa_queue_cas_write_index_scacquire_fn = HSA::hsa_queue_cas_write_index_scacquire; core_api.hsa_queue_cas_write_index_relaxed_fn = HSA::hsa_queue_cas_write_index_relaxed; - core_api.hsa_queue_cas_write_index_screlease_fn = - HSA::hsa_queue_cas_write_index_screlease; - core_api.hsa_queue_add_write_index_scacq_screl_fn = - HSA::hsa_queue_add_write_index_scacq_screl; - core_api.hsa_queue_add_write_index_scacquire_fn = - HSA::hsa_queue_add_write_index_scacquire; + core_api.hsa_queue_cas_write_index_screlease_fn = HSA::hsa_queue_cas_write_index_screlease; + core_api.hsa_queue_add_write_index_scacq_screl_fn = HSA::hsa_queue_add_write_index_scacq_screl; + core_api.hsa_queue_add_write_index_scacquire_fn = HSA::hsa_queue_add_write_index_scacquire; core_api.hsa_queue_add_write_index_relaxed_fn = HSA::hsa_queue_add_write_index_relaxed; - core_api.hsa_queue_add_write_index_screlease_fn = - HSA::hsa_queue_add_write_index_screlease; + core_api.hsa_queue_add_write_index_screlease_fn = HSA::hsa_queue_add_write_index_screlease; core_api.hsa_queue_store_read_index_relaxed_fn = HSA::hsa_queue_store_read_index_relaxed; - core_api.hsa_queue_store_read_index_screlease_fn = - HSA::hsa_queue_store_read_index_screlease; + core_api.hsa_queue_store_read_index_screlease_fn = HSA::hsa_queue_store_read_index_screlease; 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; @@ -278,14 +268,14 @@ void HsaApiTable::UpdateCore() { core_api.hsa_status_string_fn = HSA::hsa_status_string; - //Start HSA v1.1 additions + // Start HSA v1.1 additions core_api.hsa_extension_get_name_fn = HSA::hsa_extension_get_name; core_api.hsa_system_major_extension_supported_fn = HSA::hsa_system_major_extension_supported; core_api.hsa_system_get_major_extension_table_fn = HSA::hsa_system_get_major_extension_table; core_api.hsa_agent_major_extension_supported_fn = HSA::hsa_agent_major_extension_supported; core_api.hsa_cache_get_info_fn = HSA::hsa_cache_get_info; core_api.hsa_agent_iterate_caches_fn = HSA::hsa_agent_iterate_caches; - //Silent store optimization is present in all signal ops when no agents are sleeping. + // Silent store optimization is present in all signal ops when no agents are sleeping. core_api.hsa_signal_silent_store_relaxed_fn = HSA::hsa_signal_store_relaxed; core_api.hsa_signal_silent_store_screlease_fn = HSA::hsa_signal_store_screlease; core_api.hsa_signal_group_create_fn = HSA::hsa_signal_group_create; 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 9adf5476bc..be5cc744c6 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -529,24 +529,28 @@ hsa_status_t Runtime::GetSystemInfo(hsa_system_info_t attribute, void* value) { *((hsa_machine_model_t*)value) = HSA_MACHINE_MODEL_SMALL; #endif break; - case HSA_SYSTEM_INFO_EXTENSIONS: - { - memset(value, 0, sizeof(uint8_t) * 128); + case HSA_SYSTEM_INFO_EXTENSIONS: { + memset(value, 0, sizeof(uint8_t) * 128); - auto setFlag = [&](uint32_t bit) { assert(bit<128*8 && "Extension value exceeds extension bitmask"); uint index = bit/8; uint subBit = bit%8; ((uint8_t*)value)[index] |= 1<(const_cast(hsa_signals)); @@ -182,16 +181,15 @@ uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals, } } -SignalGroup::SignalGroup(uint32_t num_signals, const hsa_signal_t* hsa_signals) : count(num_signals) -{ - if(count!=0) - signals=new hsa_signal_t[count]; - else - signals=NULL; - if(signals==NULL) - return; - for(int i=0; i