From 550547f4045079df80d7a9bce9dff66ba0367e42 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Wed, 5 Jul 2017 19:33:37 -0500 Subject: [PATCH] Initial exception support for signals. Also separate signal ABI block allocations from the runtime interface object. Change-Id: If16763338db664f29163a1348f8f4c38cf0597b2 [ROCm/ROCR-Runtime commit: 2732b18092a71b729a1dc2ab5dbbc4abf35849a0] --- .../runtime/hsa-runtime/core/common/shared.h | 4 - .../hsa-runtime/core/inc/amd_aql_queue.h | 8 +- .../runtime/hsa-runtime/core/inc/checked.h | 54 ++- .../hsa-runtime/core/inc/default_signal.h | 8 +- .../hsa-runtime/core/inc/interrupt_signal.h | 8 +- .../runtime/hsa-runtime/core/inc/signal.h | 92 +++-- .../core/runtime/amd_aql_queue.cpp | 15 +- .../core/runtime/default_signal.cpp | 2 +- .../runtime/hsa-runtime/core/runtime/hsa.cpp | 356 +++++++++++++++--- .../hsa-runtime/core/runtime/hsa_ext_amd.cpp | 183 ++++++--- .../core/runtime/interrupt_signal.cpp | 5 +- 11 files changed, 543 insertions(+), 192 deletions(-) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/common/shared.h b/projects/rocr-runtime/runtime/hsa-runtime/core/common/shared.h index d56b63bcf8..b8aced5e56 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/common/shared.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/common/shared.h @@ -80,10 +80,6 @@ class Shared : public BaseShared { shared_object_ = reinterpret_cast(allocate_(sizeof(T), Max(__alignof(T), Align), 0)); - assert(shared_object_ != NULL && "Failed on allocating shared_object_"); - - memset(shared_object_, 0, sizeof(T)); - if (shared_object_ != NULL) new (shared_object_) T; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h index 1932eaa90e..08e65866c8 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -53,7 +53,7 @@ namespace amd { /// @brief Encapsulates HW Aql Command Processor functionality. It /// provide the interface for things such as Doorbell register, read, /// write pointers and a buffer. -class AqlQueue : public core::Queue, public core::Signal { +class AqlQueue : public core::Queue, private core::LocalSignal, public core::Signal { public: static __forceinline bool IsType(core::Signal* signal) { return signal->IsType(&rtti_id_); @@ -336,12 +336,6 @@ class AqlQueue : public core::Queue, public core::Signal { return NULL; } - // 64 byte-aligned allocation and release, for Queue::amd_queue_. - void* operator new(size_t size); - void* operator new(size_t size, void* ptr) { return ptr; } - void operator delete(void* ptr); - void operator delete(void*, void*) {} - protected: bool _IsA(rtti_t id) const override { return id == &rtti_id_; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/checked.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/checked.h index 5b8ef8b3c9..856d22ba8b 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/checked.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/checked.h @@ -48,29 +48,61 @@ namespace core { -/// @brief Base class for all classes whose validity can be checked using -/// IsValid() method. -template -class Checked { +/// @brief Compares type codes and pointers to check object validity. Used for cast validation. +template class Check final { public: - typedef Checked CheckedType; + typedef Check CheckType; - Checked() { object_ = uintptr_t(this) ^ uintptr_t(code); } - Checked(const Checked&) { object_ = uintptr_t(this) ^ uintptr_t(code); } - Checked(Checked&&) { object_ = uintptr_t(this) ^ uintptr_t(code); } + Check() { object_ = uintptr_t(this) ^ uintptr_t(code); } + Check(const Check&) { object_ = uintptr_t(this) ^ uintptr_t(code); } + Check(Check&&) { object_ = uintptr_t(this) ^ uintptr_t(code); } - virtual ~Checked() { object_ = NULL; } + ~Check() { object_ = NULL; } - const Checked& operator=(Checked&& rhs) { return *this; } - const Checked& operator=(const Checked& rhs) { return *this; } + const Check& operator=(Check&& rhs) { return *this; } + const Check& operator=(const Check& rhs) { return *this; } bool IsValid() const { return object_ == (uintptr_t(this) ^ uintptr_t(code)); } + uint64_t check_code() const { return code; } + private: uintptr_t object_; }; +template class Check final { + public: + typedef Check CheckType; + + Check() { object_ = uintptr_t(code); } + Check(const Check&) { object_ = uintptr_t(code); } + Check(Check&&) { object_ = uintptr_t(code); } + + const Check& operator=(Check&& rhs) { return *this; } + const Check& operator=(const Check& rhs) { return *this; } + + bool IsValid() const { return object_ == uintptr_t(code); } + + uint64_t check_code() const { return code; } + + private: + uintptr_t object_; +}; + +/// @brief Base class for validating objects. +template class Checked { + public: + typedef Checked CheckedType; + + bool IsValid() const { return id.IsValid(); } + + virtual ~Checked() {} + + private: + Check id; +}; + } // namespace core #endif // header guard diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/default_signal.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/default_signal.h index f0f13eb06e..6f3377b28c 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/default_signal.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/default_signal.h @@ -53,7 +53,7 @@ namespace core { /// @brief Simple pure memory based signal. /// @brief See base class Signal. -class DefaultSignal : public Signal { +class DefaultSignal : private LocalSignal, public Signal { public: /// @brief Determines if a Signal* can be safely converted to DefaultSignal* /// via static_cast. @@ -155,12 +155,6 @@ class DefaultSignal : public Signal { /// @brief see the base class Signal __forceinline HsaEvent* EopEvent() { return NULL; } - /// @brief prevent throwing exceptions - void* operator new(size_t size) { return malloc(size); } - - /// @brief prevent throwing exceptions - void operator delete(void* ptr) { free(ptr); } - protected: bool _IsA(rtti_t id) const { return id == &rtti_id_; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/interrupt_signal.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/interrupt_signal.h index adbbb5070b..755f0dd447 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/interrupt_signal.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/interrupt_signal.h @@ -61,7 +61,7 @@ namespace core { /// Supports only one waiter for now. /// KFD changes are needed to support multiple waiters and have device /// signaling. -class InterruptSignal : public Signal { +class InterruptSignal : private LocalSignal, public Signal { public: static HsaEvent* CreateEvent(HSA_EVENTTYPE type, bool manual_reset); static void DestroyEvent(HsaEvent* evt); @@ -165,12 +165,6 @@ class InterruptSignal : public Signal { /// @brief See base class Signal. __forceinline HsaEvent* EopEvent() { return event_; } - /// @brief prevent throwing exceptions - void* operator new(size_t size) { return malloc(size); } - - /// @brief prevent throwing exceptions - void operator delete(void* ptr) { free(ptr); } - protected: bool _IsA(rtti_t id) const { return id == &rtti_id_; } 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 21e61bdcae..502a7880c0 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/signal.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/signal.h @@ -53,6 +53,7 @@ #include "core/inc/checked.h" #include "core/util/utils.h" +#include "core/util/locks.h" #include "inc/amd_hsa_signal.h" @@ -60,54 +61,72 @@ namespace core { class Agent; class Signal; -/// @brief Helper structure to simplify conversion of amd_signal_t and -/// core::Signal object. +/// @brief ABI and object conversion struct for signals. May be shared between processes. struct SharedSignal { amd_signal_t amd_signal; Signal* core_signal; + Check<0x71FCCA6A3D5D5276, true> id; + + bool IsValid() const { return id.IsValid(); } + + SharedSignal() { + amd_signal.kind = AMD_SIGNAL_KIND_INVALID; + core_signal = nullptr; + } + + static __forceinline SharedSignal* Convert(hsa_signal_t signal) { + if (signal.handle == 0) throw std::bad_cast(); + SharedSignal* ret = reinterpret_cast(static_cast(signal.handle) - + offsetof(SharedSignal, amd_signal)); + if (!ret->IsValid()) throw std::bad_cast(); + return ret; + } +}; +static_assert(std::is_standard_layout::value, + "SharedSignal must remain standard layout for IPC use."); +static_assert(std::is_trivially_destructible::value, + "SharedSignal must not be modified on delete for IPC use."); + +class LocalSignal { + public: + explicit LocalSignal(hsa_signal_value_t initial_value) { + if (!local_signal_.IsSharedObjectAllocationValid()) throw std::bad_alloc(); + local_signal_.shared_object()->amd_signal.value = initial_value; + } + + SharedSignal* signal() const { return local_signal_.shared_object(); } + + private: + Shared local_signal_; }; /// @brief An abstract base class which helps implement the public hsa_signal_t /// type (an opaque handle) and its associated APIs. At its core, signal uses /// a 32 or 64 bit value. This value can be waitied on or signaled atomically /// using specified memory ordering semantics. -class Signal : public Checked<0x71FCCA6A3D5D5276>, - public Shared { +class Signal { public: - /// @brief Constructor initializes the signal with initial value. - explicit Signal(hsa_signal_value_t initial_value) - : Shared(), - signal_(shared_object()->amd_signal), - async_copy_agent_(NULL) { - if (!Shared::IsSharedObjectAllocationValid()) { - invalid_ = true; - return; - } + /// @brief Constructor Links and publishes the signal interface object. + explicit Signal(SharedSignal* abi_block) + : signal_(abi_block->amd_signal), async_copy_agent_(NULL) { + if (abi_block == nullptr) throw std::bad_alloc(); - shared_object()->core_signal = this; - - signal_.kind = AMD_SIGNAL_KIND_INVALID; - signal_.value = initial_value; invalid_ = false; waiting_ = 0; retained_ = 0; + + abi_block->core_signal = this; } - virtual ~Signal() { signal_.kind = AMD_SIGNAL_KIND_INVALID; } + virtual ~Signal() { invalid_ = true; } - bool IsValid() const { - if (CheckedType::IsValid() && !invalid_) return true; - return false; - } + bool IsValid() const { return !invalid_; } /// @brief Converts from this implementation class to the public /// hsa_signal_t type - an opaque handle. static __forceinline hsa_signal_t Convert(Signal* signal) { - const uint64_t handle = - (signal != NULL && signal->IsValid()) - ? static_cast( - reinterpret_cast(&signal->signal_)) - : 0; + assert(signal != nullptr && signal->IsValid() && "Conversion on invalid Signal object."); + const uint64_t handle = static_cast(reinterpret_cast(&signal->signal_)); const hsa_signal_t signal_handle = {handle}; return signal_handle; } @@ -115,11 +134,8 @@ class Signal : public Checked<0x71FCCA6A3D5D5276>, /// @brief Converts from this implementation class to the public /// hsa_signal_t type - an opaque handle. static __forceinline const hsa_signal_t Convert(const Signal* signal) { - const uint64_t handle = - (signal != NULL && signal->IsValid()) - ? static_cast( - reinterpret_cast(&signal->signal_)) - : 0; + assert(signal != nullptr && signal->IsValid() && "Conversion on invalid Signal object."); + const uint64_t handle = static_cast(reinterpret_cast(&signal->signal_)); const hsa_signal_t signal_handle = {handle}; return signal_handle; } @@ -127,13 +143,11 @@ class Signal : public Checked<0x71FCCA6A3D5D5276>, /// @brief Converts from public hsa_signal_t type (an opaque handle) to /// this implementation class object. static __forceinline Signal* Convert(hsa_signal_t signal) { - return (signal.handle != 0) - ? reinterpret_cast( - static_cast(signal.handle) - - (reinterpret_cast( - &reinterpret_cast(1234)->amd_signal) - - uintptr_t(1234)))->core_signal - : NULL; + SharedSignal* shared = SharedSignal::Convert(signal); + if (shared->core_signal != nullptr) + return shared->core_signal; + else + throw std::bad_cast(); } // Below are various methods corresponding to the APIs, which load/store the 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 9d680ace43..b79a1b249d 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 @@ -77,18 +77,11 @@ volatile uint32_t AqlQueue::queue_count_ = 0; KernelMutex AqlQueue::queue_lock_; int AqlQueue::rtti_id_; -void* AqlQueue::operator new(size_t size) { - // Align base to 64B to enforce amd_queue_ member alignment. - return _aligned_malloc(size, kAmdQueueAlignBytes); -} - -void AqlQueue::operator delete(void* ptr) { _aligned_free(ptr); } - -AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, - ScratchInfo& scratch, core::HsaEventCallback callback, - void* err_data, bool is_kv) +AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, ScratchInfo& scratch, + core::HsaEventCallback callback, void* err_data, bool is_kv) : Queue(), - Signal(0), + LocalSignal(0), + Signal(signal()), ring_buf_(NULL), ring_buf_alloc_bytes_(0), queue_id_(HSA_QUEUEID(-1)), diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp index 44348d03c0..b1a89a5af7 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp @@ -48,7 +48,7 @@ namespace core { int DefaultSignal::rtti_id_ = 0; DefaultSignal::DefaultSignal(hsa_signal_value_t initial_value) - : Signal(initial_value) { + : LocalSignal(initial_value), Signal(signal()) { signal_.kind = AMD_SIGNAL_KIND_USER; signal_.event_mailbox_ptr = NULL; HSA::hsa_memory_register(this, sizeof(DefaultSignal)); 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 04217e21c5..890f642852 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -163,6 +163,20 @@ static __forceinline bool IsValid(T* ptr) { return (ptr == NULL) ? NULL : ptr->IsValid(); } +namespace AMD { +hsa_status_t handleException(); + +template static __forceinline T handleExceptionT() { + handleException(); + abort(); + return T(); +} +} + +#define TRY try { +#define CATCH } catch(...) { return AMD::handleException(); } +#define CATCHRET(RETURN_TYPE) } catch(...) { return AMD::handleExceptionT(); } + //----------------------------------------------------------------------------- // Basic Checks //----------------------------------------------------------------------------- @@ -184,7 +198,7 @@ static_assert(sizeof(void*) == 4, "HSA_LARGE_MODEL is set incorrectly!"); #endif #if !defined(HSA_LARGE_MODEL) || !defined(__linux__) -static_assert(false, "Only HSA_LARGE_MODEL (64bit mode) and Linux supported."); +// static_assert(false, "Only HSA_LARGE_MODEL (64bit mode) and Linux supported."); #endif namespace HSA { @@ -192,11 +206,17 @@ namespace HSA { //---------------------------------------------------------------------------// // Init/Shutdown routines //---------------------------------------------------------------------------// -hsa_status_t hsa_init() { return core::Runtime::runtime_singleton_->Acquire(); } +hsa_status_t hsa_init() { + TRY + return core::Runtime::runtime_singleton_->Acquire(); + CATCH +} hsa_status_t hsa_shut_down() { + TRY IS_OPEN(); return core::Runtime::runtime_singleton_->Release(); + CATCH } //---------------------------------------------------------------------------// @@ -204,12 +224,15 @@ hsa_status_t hsa_shut_down() { //---------------------------------------------------------------------------// hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void* value) { + TRY IS_OPEN(); IS_BAD_PTR(value); return core::Runtime::runtime_singleton_->GetSystemInfo(attribute, value); + CATCH } hsa_status_t hsa_extension_get_name(uint16_t extension, const char** name) { + TRY IS_OPEN(); IS_BAD_PTR(name); switch (extension) { @@ -239,11 +262,13 @@ hsa_status_t hsa_extension_get_name(uint16_t extension, const char** name) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_system_extension_supported(uint16_t extension, uint16_t version_major, uint16_t version_minor, bool* result) { + TRY IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && @@ -271,10 +296,12 @@ hsa_status_t } return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_system_major_extension_supported(uint16_t extension, uint16_t version_major, uint16_t* version_minor, bool* result) { + TRY IS_OPEN(); IS_BAD_PTR(version_minor); IS_BAD_PTR(result); @@ -305,6 +332,7 @@ hsa_status_t hsa_system_major_extension_supported(uint16_t extension, uint16_t v *result = false; return HSA_STATUS_SUCCESS; + CATCH } static size_t get_extension_table_length(uint16_t extension, uint16_t major, uint16_t minor) { @@ -362,13 +390,16 @@ static size_t get_extension_table_length(uint16_t extension, uint16_t major, uin hsa_status_t hsa_system_get_extension_table(uint16_t extension, uint16_t version_major, uint16_t version_minor, void* table) { + TRY return HSA::hsa_system_get_major_extension_table( extension, version_major, get_extension_table_length(extension, version_major, version_minor), table); + CATCH } hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t version_major, size_t table_length, void* table) { + TRY IS_OPEN(); IS_BAD_PTR(table); @@ -453,6 +484,7 @@ hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t v } return HSA_STATUS_ERROR; + CATCH } //---------------------------------------------------------------------------// @@ -461,24 +493,29 @@ hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t v hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, void* data), void* data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); return core::Runtime::runtime_singleton_->IterateAgent(callback, data); + CATCH } hsa_status_t hsa_agent_get_info(hsa_agent_t agent_handle, hsa_agent_info_t attribute, void* value) { + TRY IS_OPEN(); IS_BAD_PTR(value); const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); return agent->GetInfo(attribute, value); + CATCH } hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent_handle, hsa_profile_t profile, uint16_t* mask) { + TRY IS_OPEN(); IS_BAD_PTR(mask); IS_BAD_PROFILE(profile); @@ -487,30 +524,36 @@ hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent_handle, *mask = 0; return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, void* value) { + TRY IS_OPEN(); core::Cache* Cache = core::Cache::Convert(cache); IS_VALID(Cache); IS_BAD_PTR(value); return Cache->GetInfo(attribute, value); + CATCH } hsa_status_t hsa_agent_iterate_caches(hsa_agent_t agent_handle, hsa_status_t (*callback)(hsa_cache_t cache, void* data), void* data) { + TRY IS_OPEN(); const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); IS_BAD_PTR(callback); return agent->IterateCache(callback, data); + CATCH } hsa_status_t hsa_agent_extension_supported(uint16_t extension, hsa_agent_t agent_handle, uint16_t version_major, uint16_t version_minor, bool* result) { + TRY IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && @@ -539,11 +582,13 @@ hsa_status_t } return HSA_STATUS_SUCCESS; + CATCH } 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) { + TRY IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && @@ -569,6 +614,7 @@ hsa_status_t hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t } return HSA_STATUS_SUCCESS; + CATCH } /// @brief Api to create a user mode queue. @@ -593,6 +639,7 @@ hsa_status_t hsa_queue_create( void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data, uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue) { + TRY IS_OPEN(); if ((queue == NULL) || (size == 0) || (!IsPowerOfTwo(size)) || @@ -629,12 +676,14 @@ hsa_status_t hsa_queue_create( } return status; + CATCH } hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size, hsa_queue_type32_t type, uint32_t features, hsa_signal_t doorbell_signal, hsa_queue_t** queue) { + TRY IS_OPEN(); if ((queue == NULL) || (region.handle == 0) || @@ -661,6 +710,7 @@ hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size, *queue = core::Queue::Convert(host_queue); return HSA_STATUS_SUCCESS; + CATCH } /// @brief Api to destroy a user mode queue @@ -669,12 +719,14 @@ hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size, /// /// @return hsa_status hsa_status_t hsa_queue_destroy(hsa_queue_t* queue) { + TRY IS_OPEN(); IS_BAD_PTR(queue); core::Queue* cmd_queue = core::Queue::Convert(queue); IS_VALID(cmd_queue); delete cmd_queue; return HSA_STATUS_SUCCESS; + CATCH } /// @brief Api to inactivate a user mode queue @@ -683,12 +735,14 @@ hsa_status_t hsa_queue_destroy(hsa_queue_t* queue) { /// /// @return hsa_status hsa_status_t hsa_queue_inactivate(hsa_queue_t* queue) { + TRY IS_OPEN(); IS_BAD_PTR(queue); core::Queue* cmd_queue = core::Queue::Convert(queue); IS_VALID(cmd_queue); cmd_queue->Inactivate(); return HSA_STATUS_SUCCESS; + CATCH } /// @brief Api to read the Read Index of Queue using Acquire semantics @@ -697,9 +751,11 @@ hsa_status_t hsa_queue_inactivate(hsa_queue_t* queue) { /// /// @return uint64_t Value of Read index uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t* queue) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadReadIndexAcquire(); + CATCHRET(uint64_t) } /// @brief Api to read the Read Index of Queue using Relaxed semantics @@ -708,9 +764,11 @@ uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t* queue) { /// /// @return uint64_t Value of Read index uint64_t hsa_queue_load_read_index_relaxed(const hsa_queue_t* queue) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadReadIndexRelaxed(); + CATCHRET(uint64_t) } /// @brief Api to read the Write Index of Queue using Acquire semantics @@ -719,9 +777,11 @@ uint64_t hsa_queue_load_read_index_relaxed(const hsa_queue_t* queue) { /// /// @return uint64_t Value of Write index uint64_t hsa_queue_load_write_index_scacquire(const hsa_queue_t* queue) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadWriteIndexAcquire(); + CATCHRET(uint64_t) } /// @brief Api to read the Write Index of Queue using Relaxed semantics @@ -730,9 +790,11 @@ uint64_t hsa_queue_load_write_index_scacquire(const hsa_queue_t* queue) { /// /// @return uint64_t Value of Write index uint64_t hsa_queue_load_write_index_relaxed(const hsa_queue_t* queue) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadWriteIndexAcquire(); + CATCHRET(uint64_t) } /// @brief Api to store the Read Index of Queue using Relaxed semantics @@ -742,9 +804,11 @@ uint64_t hsa_queue_load_write_index_relaxed(const hsa_queue_t* queue) { /// @param value Value of new read index void hsa_queue_store_read_index_relaxed(const hsa_queue_t* queue, uint64_t value) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreReadIndexRelaxed(value); + CATCHRET(void) } /// @brief Api to store the Read Index of Queue using Release semantics @@ -753,9 +817,11 @@ void hsa_queue_store_read_index_relaxed(const hsa_queue_t* queue, /// /// @param value Value of new read index void hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, uint64_t value) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreReadIndexRelease(value); + CATCHRET(void) } /// @brief Api to store the Write Index of Queue using Relaxed semantics @@ -765,9 +831,11 @@ void hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, uint64_t val /// @param value Value of new write index void hsa_queue_store_write_index_relaxed(const hsa_queue_t* queue, uint64_t value) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreWriteIndexRelaxed(value); + CATCHRET(void) } /// @brief Api to store the Write Index of Queue using Release semantics @@ -776,9 +844,11 @@ void hsa_queue_store_write_index_relaxed(const hsa_queue_t* queue, /// /// @param value Value of new write index void hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreWriteIndexRelease(value); + CATCHRET(void) } /// @brief Api to compare and swap the Write Index of Queue using Acquire and @@ -793,9 +863,11 @@ void hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, uint64_t va /// @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) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcqRel(expected, value); + CATCHRET(uint64_t) } /// @brief Api to compare and swap the Write Index of Queue using Acquire @@ -810,9 +882,11 @@ uint64_t hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, uint64_ /// @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) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcquire(expected, value); + CATCHRET(uint64_t) } /// @brief Api to compare and swap the Write Index of Queue using Relaxed @@ -828,9 +902,11 @@ uint64_t hsa_queue_cas_write_index_scacquire(const hsa_queue_t* queue, uint64_t uint64_t hsa_queue_cas_write_index_relaxed(const hsa_queue_t* queue, uint64_t expected, uint64_t value) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexRelaxed(expected, value); + CATCHRET(uint64_t) } /// @brief Api to compare and swap the Write Index of Queue using Release @@ -845,9 +921,11 @@ uint64_t hsa_queue_cas_write_index_relaxed(const hsa_queue_t* queue, /// @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) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexRelease(expected, value); + CATCHRET(uint64_t) } /// @brief Api to Add to the Write Index of Queue using Acquire and Release @@ -859,9 +937,11 @@ uint64_t hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, uint64_t /// /// @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) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcqRel(value); + CATCHRET(uint64_t) } /// @brief Api to Add to the Write Index of Queue using Acquire Semantics @@ -872,9 +952,11 @@ uint64_t hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, uint64_ /// /// @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) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcquire(value); + CATCHRET(uint64_t) } /// @brief Api to Add to the Write Index of Queue using Relaxed Semantics @@ -886,9 +968,11 @@ uint64_t hsa_queue_add_write_index_scacquire(const hsa_queue_t* queue, uint64_t /// @return uint64_t Value of write index before the update uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t* queue, uint64_t value) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexRelaxed(value); + CATCHRET(uint64_t) } /// @brief Api to Add to the Write Index of Queue using Release Semantics @@ -899,9 +983,11 @@ uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t* queue, /// /// @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) { + TRY core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexRelease(value); + CATCHRET(uint64_t) } //----------------------------------------------------------------------------- @@ -910,16 +996,19 @@ uint64_t hsa_queue_add_write_index_screlease(const hsa_queue_t* queue, uint64_t hsa_status_t hsa_agent_iterate_regions( hsa_agent_t agent_handle, hsa_status_t (*callback)(hsa_region_t region, void* data), void* data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); return agent->IterateRegion(callback, data); + CATCH } hsa_status_t hsa_region_get_info(hsa_region_t region, hsa_region_info_t attribute, void* value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -927,9 +1016,11 @@ hsa_status_t hsa_region_get_info(hsa_region_t region, IS_VALID(mem_region); return mem_region->GetInfo(attribute, value); + CATCH } hsa_status_t hsa_memory_register(void* address, size_t size) { + TRY IS_OPEN(); if (size == 0 && address != NULL) { @@ -937,16 +1028,20 @@ hsa_status_t hsa_memory_register(void* address, size_t size) { } return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_memory_deregister(void* address, size_t size) { + TRY IS_OPEN(); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_memory_allocate(hsa_region_t region, size_t size, void** ptr) { + TRY IS_OPEN(); if (size == 0 || ptr == NULL) { @@ -958,9 +1053,11 @@ hsa_status_t return core::Runtime::runtime_singleton_->AllocateMemory( mem_region, size, core::MemoryRegion::AllocateNoFlags, ptr); + CATCH } hsa_status_t hsa_memory_free(void* ptr) { + TRY IS_OPEN(); if (ptr == NULL) { @@ -968,11 +1065,13 @@ hsa_status_t hsa_memory_free(void* ptr) { } return core::Runtime::runtime_singleton_->FreeMemory(ptr); + CATCH } hsa_status_t hsa_memory_assign_agent(void* ptr, hsa_agent_t agent_handle, hsa_access_permission_t access) { + TRY IS_OPEN(); if ((ptr == NULL) || (access < HSA_ACCESS_PERMISSION_RO) || @@ -984,9 +1083,11 @@ hsa_status_t hsa_memory_assign_agent(void* ptr, IS_VALID(agent); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_memory_copy(void* dst, const void* src, size_t size) { + TRY IS_OPEN(); if (dst == NULL || src == NULL) { @@ -998,21 +1099,23 @@ hsa_status_t hsa_memory_copy(void* dst, const void* src, size_t size) { } return core::Runtime::runtime_singleton_->CopyMemory(dst, src, size); + CATCH } //----------------------------------------------------------------------------- // Signals //----------------------------------------------------------------------------- -typedef struct { +struct AgentHandleCompare { bool operator()(const hsa_agent_t& lhs, const hsa_agent_t& rhs) const { return lhs.handle < rhs.handle; } -} AgentHandleCompare; +}; hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, const hsa_agent_t* consumers, hsa_signal_t* hsa_signal) { + TRY IS_OPEN(); IS_BAD_PTR(hsa_signal); @@ -1055,9 +1158,11 @@ hsa_status_t } return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_signal_destroy(hsa_signal_t hsa_signal) { + TRY IS_OPEN(); if (hsa_signal.handle == 0) { @@ -1068,31 +1173,40 @@ hsa_status_t hsa_signal_destroy(hsa_signal_t hsa_signal) { IS_VALID(signal); delete signal; return HSA_STATUS_SUCCESS; + CATCH } hsa_signal_value_t hsa_signal_load_relaxed(hsa_signal_t hsa_signal) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->LoadRelaxed(); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t hsa_signal) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->LoadAcquire(); + CATCHRET(hsa_signal_value_t) } void hsa_signal_store_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->StoreRelaxed(value); + CATCHRET(void) } void hsa_signal_store_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->StoreRelease(value); + CATCHRET(void) } hsa_signal_value_t @@ -1101,10 +1215,12 @@ hsa_signal_value_t hsa_signal_value_t compare_value, uint64_t timeout_hint, hsa_wait_state_t wait_state_hint) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->WaitRelaxed(condition, compare_value, timeout_hint, wait_state_hint); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t hsa_signal, @@ -1112,15 +1228,18 @@ hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t compare_value, uint64_t timeout_hint, hsa_wait_state_t wait_state_hint) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->WaitAcquire(condition, compare_value, timeout_hint, wait_state_hint); + CATCHRET(hsa_signal_value_t) } 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) { + TRY 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])); @@ -1133,14 +1252,17 @@ hsa_status_t hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t* s } *signal_group = core::SignalGroup::Convert(group); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_signal_group_destroy(hsa_signal_group_t signal_group) { + TRY IS_OPEN(); core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); delete group; return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, @@ -1148,6 +1270,7 @@ hsa_status_t hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, const hsa_signal_value_t* compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t* signal, hsa_signal_value_t* value) { + TRY IS_OPEN(); const core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); @@ -1158,6 +1281,7 @@ hsa_status_t hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, if (index >= group->Count()) return HSA_STATUS_ERROR_INVALID_ARGUMENT; *signal = group->List()[index]; return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, @@ -1165,193 +1289,247 @@ hsa_status_t hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group const hsa_signal_value_t* compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t* signal, hsa_signal_value_t* value) { + TRY 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; + CATCH } -void - hsa_signal_and_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_and_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndRelaxed(value); + CATCHRET(void) } void hsa_signal_and_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcquire(value); + CATCHRET(void) } void hsa_signal_and_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndRelease(value); + CATCHRET(void) } void hsa_signal_and_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcqRel(value); + CATCHRET(void) } -void - hsa_signal_or_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_or_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrRelaxed(value); + CATCHRET(void) } void hsa_signal_or_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcquire(value); + CATCHRET(void) } void hsa_signal_or_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrRelease(value); + CATCHRET(void) } void hsa_signal_or_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcqRel(value); + CATCHRET(void) } -void - hsa_signal_xor_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_xor_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorRelaxed(value); + CATCHRET(void) } void hsa_signal_xor_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcquire(value); + CATCHRET(void) } void hsa_signal_xor_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorRelease(value); + CATCHRET(void) } void hsa_signal_xor_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcqRel(value); + CATCHRET(void) } -void - hsa_signal_add_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { +void hsa_signal_add_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->AddRelaxed(value); + CATCHRET(void) } void hsa_signal_add_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcquire(value); + CATCHRET(void) } void hsa_signal_add_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddRelease(value); + CATCHRET(void) } void hsa_signal_add_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcqRel(value); + CATCHRET(void) } void hsa_signal_subtract_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubRelaxed(value); + CATCHRET(void) } void hsa_signal_subtract_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcquire(value); + CATCHRET(void) } void hsa_signal_subtract_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubRelease(value); + CATCHRET(void) } void hsa_signal_subtract_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcqRel(value); + CATCHRET(void) } hsa_signal_value_t hsa_signal_exchange_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchRelaxed(value); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_exchange_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcquire(value); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_exchange_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchRelease(value); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_exchange_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcqRel(value); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_cas_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasRelaxed(expected, value); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_cas_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcquire(expected, value); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_cas_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasRelease(expected, value); + CATCHRET(hsa_signal_value_t) } hsa_signal_value_t hsa_signal_cas_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { + TRY core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcqRel(expected, value); + CATCHRET(hsa_signal_value_t) } //===--- Instruction Set Architecture -------------------------------------===// @@ -1363,6 +1541,7 @@ using core::Wavefront; hsa_status_t hsa_isa_from_name( const char *name, hsa_isa_t *isa) { + TRY IS_OPEN(); IS_BAD_PTR(name); IS_BAD_PTR(isa); @@ -1374,6 +1553,7 @@ hsa_status_t hsa_isa_from_name( *isa = Isa::Handle(isa_object); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_agent_iterate_isas( @@ -1381,6 +1561,7 @@ hsa_status_t hsa_agent_iterate_isas( hsa_status_t (*callback)(hsa_isa_t isa, void *data), void *data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); @@ -1393,13 +1574,16 @@ hsa_status_t hsa_agent_iterate_isas( } return callback(Isa::Handle(isa_object), data); + CATCH } -/* deprecated */ hsa_status_t hsa_isa_get_info( +/* deprecated */ +hsa_status_t hsa_isa_get_info( hsa_isa_t isa, hsa_isa_info_t attribute, uint32_t index, void *value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -1412,12 +1596,14 @@ hsa_status_t hsa_agent_iterate_isas( return isa_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; + CATCH } hsa_status_t hsa_isa_get_info_alt( hsa_isa_t isa, hsa_isa_info_t attribute, void *value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -1426,12 +1612,14 @@ hsa_status_t hsa_isa_get_info_alt( return isa_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; + CATCH } hsa_status_t hsa_isa_get_exception_policies( hsa_isa_t isa, hsa_profile_t profile, uint16_t *mask) { + TRY IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_PTR(mask); @@ -1442,6 +1630,7 @@ hsa_status_t hsa_isa_get_exception_policies( // FIXME: update when exception policies are supported. *mask = 0; return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_isa_get_round_method( @@ -1449,6 +1638,7 @@ hsa_status_t hsa_isa_get_round_method( hsa_fp_type_t fp_type, hsa_flush_mode_t flush_mode, hsa_round_method_t *round_method) { + TRY IS_OPEN(); IS_BAD_FP_TYPE(fp_type); IS_BAD_FLUSH_MODE(flush_mode); @@ -1459,12 +1649,14 @@ hsa_status_t hsa_isa_get_round_method( *round_method = isa_object->GetRoundMethod(fp_type, flush_mode); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_wavefront_get_info( hsa_wavefront_t wavefront, hsa_wavefront_info_t attribute, void *value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -1475,6 +1667,7 @@ hsa_status_t hsa_wavefront_get_info( return wavefront_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; + CATCH } hsa_status_t hsa_isa_iterate_wavefronts( @@ -1482,6 +1675,7 @@ hsa_status_t hsa_isa_iterate_wavefronts( hsa_status_t (*callback)(hsa_wavefront_t wavefront, void *data), void *data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); @@ -1492,12 +1686,15 @@ hsa_status_t hsa_isa_iterate_wavefronts( assert(wavefront_object); return callback(Wavefront::Handle(wavefront_object), data); + CATCH } -/* deprecated */ hsa_status_t hsa_isa_compatible( +/* deprecated */ +hsa_status_t hsa_isa_compatible( hsa_isa_t code_object_isa, hsa_isa_t agent_isa, bool *result) { + TRY IS_OPEN(); IS_BAD_PTR(result); @@ -1509,6 +1706,7 @@ hsa_status_t hsa_isa_iterate_wavefronts( *result = code_object_isa_object->IsCompatible(agent_isa_object); return HSA_STATUS_SUCCESS; + CATCH } //===--- Code Objects (deprecated) ----------------------------------------===// @@ -1573,7 +1771,8 @@ AmdHsaCodeManager *GetCodeManager() { } // namespace anonymous -/* deprecated */ hsa_status_t hsa_code_object_serialize( +/* deprecated */ +hsa_status_t hsa_code_object_serialize( hsa_code_object_t code_object, hsa_status_t (*alloc_callback)(size_t size, hsa_callback_data_t data, @@ -1582,6 +1781,7 @@ AmdHsaCodeManager *GetCodeManager() { const char *options, void **serialized_code_object, size_t *serialized_code_object_size) { + TRY IS_OPEN(); IS_BAD_PTR(alloc_callback); IS_BAD_PTR(serialized_code_object); @@ -1603,13 +1803,16 @@ AmdHsaCodeManager *GetCodeManager() { *serialized_code_object_size = code->ElfSize(); return HSA_STATUS_SUCCESS; + CATCH } -/* deprecated */ hsa_status_t hsa_code_object_deserialize( +/* deprecated */ +hsa_status_t hsa_code_object_deserialize( void *serialized_code_object, size_t serialized_code_object_size, const char *options, hsa_code_object_t *code_object) { + TRY IS_OPEN(); IS_BAD_PTR(serialized_code_object); IS_BAD_PTR(code_object); @@ -1640,10 +1843,13 @@ AmdHsaCodeManager *GetCodeManager() { code_object->handle = reinterpret_cast(code_object_alloc_data); return HSA_STATUS_SUCCESS; + CATCH } -/* deprecated */ hsa_status_t hsa_code_object_destroy( +/* deprecated */ +hsa_status_t hsa_code_object_destroy( hsa_code_object_t code_object) { + TRY IS_OPEN(); void *code_object_data = reinterpret_cast(code_object.handle); @@ -1657,12 +1863,15 @@ AmdHsaCodeManager *GetCodeManager() { HSA::hsa_memory_free(code_object_data); return HSA_STATUS_SUCCESS; + CATCH } -/* deprecated */ hsa_status_t hsa_code_object_get_info( +/* deprecated */ +hsa_status_t hsa_code_object_get_info( hsa_code_object_t code_object, hsa_code_object_info_t attribute, void *value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -1672,32 +1881,35 @@ AmdHsaCodeManager *GetCodeManager() { } switch (attribute) { - case HSA_CODE_OBJECT_INFO_ISA: { - char isa_name[64]; - hsa_status_t status = code->GetInfo(attribute, &isa_name); - if (status != HSA_STATUS_SUCCESS) { - return status; + case HSA_CODE_OBJECT_INFO_ISA: { + char isa_name[64]; + hsa_status_t status = code->GetInfo(attribute, &isa_name); + if (status != HSA_STATUS_SUCCESS) { + return status; + } + + hsa_isa_t isa_handle = {0}; + status = HSA::hsa_isa_from_name(isa_name, &isa_handle); + if (status != HSA_STATUS_SUCCESS) { + return status; + } + + *((hsa_isa_t*)value) = isa_handle; + return HSA_STATUS_SUCCESS; } - - hsa_isa_t isa_handle = {0}; - status = HSA::hsa_isa_from_name(isa_name, &isa_handle); - if (status != HSA_STATUS_SUCCESS) { - return status; + default: { + return code->GetInfo(attribute, value); } - - *((hsa_isa_t*)value) = isa_handle; - return HSA_STATUS_SUCCESS; - } - default: { - return code->GetInfo(attribute, value); - } } + CATCH } -/* deprecated */ hsa_status_t hsa_code_object_get_symbol( +/* deprecated */ +hsa_status_t hsa_code_object_get_symbol( hsa_code_object_t code_object, const char *symbol_name, hsa_code_symbol_t *symbol) { + TRY IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -1708,13 +1920,16 @@ AmdHsaCodeManager *GetCodeManager() { } return code->GetSymbol(nullptr, symbol_name, symbol); + CATCH } -/* deprecated */ hsa_status_t hsa_code_object_get_symbol_from_name( +/* deprecated */ +hsa_status_t hsa_code_object_get_symbol_from_name( hsa_code_object_t code_object, const char *module_name, const char *symbol_name, hsa_code_symbol_t *symbol) { + TRY IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -1725,12 +1940,15 @@ AmdHsaCodeManager *GetCodeManager() { } return code->GetSymbol(module_name, symbol_name, symbol); + CATCH } -/* deprecated */ hsa_status_t hsa_code_symbol_get_info( +/* deprecated */ +hsa_status_t hsa_code_symbol_get_info( hsa_code_symbol_t code_symbol, hsa_code_symbol_info_t attribute, void *value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -1740,14 +1958,17 @@ AmdHsaCodeManager *GetCodeManager() { } return symbol->GetInfo(attribute, value); + CATCH } -/* deprecated */ hsa_status_t hsa_code_object_iterate_symbols( +/* deprecated */ +hsa_status_t hsa_code_object_iterate_symbols( hsa_code_object_t code_object, hsa_status_t (*callback)(hsa_code_object_t code_object, hsa_code_symbol_t symbol, void *data), void *data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); @@ -1757,6 +1978,7 @@ AmdHsaCodeManager *GetCodeManager() { } return code->IterateSymbols(code_object, callback, data); + CATCH } //===--- Executable -------------------------------------------------------===// @@ -1810,6 +2032,7 @@ Loader *GetLoader() { hsa_status_t hsa_code_object_reader_create_from_file( hsa_file_t file, hsa_code_object_reader_t *code_object_reader) { + TRY IS_OPEN(); IS_BAD_PTR(code_object_reader); @@ -1839,12 +2062,14 @@ hsa_status_t hsa_code_object_reader_create_from_file( *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_code_object_reader_create_from_memory( const void *code_object, size_t size, hsa_code_object_reader_t *code_object_reader) { + TRY IS_OPEN(); IS_BAD_PTR(code_object); IS_BAD_PTR(code_object_reader); @@ -1859,10 +2084,12 @@ hsa_status_t hsa_code_object_reader_create_from_memory( *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_code_object_reader_destroy( hsa_code_object_reader_t code_object_reader) { + TRY IS_OPEN(); CodeObjectReaderWrapper *wrapper = CodeObjectReaderWrapper::Object( @@ -1877,13 +2104,16 @@ hsa_status_t hsa_code_object_reader_destroy( delete wrapper; return HSA_STATUS_SUCCESS; + CATCH } -/* deprecated */ hsa_status_t hsa_executable_create( +/* deprecated */ +hsa_status_t hsa_executable_create( hsa_profile_t profile, hsa_executable_state_t executable_state, const char *options, hsa_executable_t *executable) { + TRY IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_EXECUTABLE_STATE(executable_state); @@ -1906,6 +2136,7 @@ hsa_status_t hsa_code_object_reader_destroy( } return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_executable_create_alt( @@ -1913,6 +2144,7 @@ hsa_status_t hsa_executable_create_alt( hsa_default_float_rounding_mode_t default_float_rounding_mode, const char *options, hsa_executable_t *executable) { + TRY IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_ROUNDING_MODE(default_float_rounding_mode); // NOTES: should we check @@ -1927,10 +2159,12 @@ hsa_status_t hsa_executable_create_alt( *executable = Executable::Handle(exec); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_executable_destroy( hsa_executable_t executable) { + TRY IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -1940,13 +2174,16 @@ hsa_status_t hsa_executable_destroy( GetLoader()->DestroyExecutable(exec); return HSA_STATUS_SUCCESS; + CATCH } -/* deprecated */ hsa_status_t hsa_executable_load_code_object( +/* deprecated */ +hsa_status_t hsa_executable_load_code_object( hsa_executable_t executable, hsa_agent_t agent, hsa_code_object_t code_object, const char *options) { + TRY IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -1955,6 +2192,7 @@ hsa_status_t hsa_executable_destroy( } return exec->LoadCodeObject(agent, code_object, options); + CATCH } hsa_status_t hsa_executable_load_program_code_object( @@ -1962,6 +2200,7 @@ hsa_status_t hsa_executable_load_program_code_object( hsa_code_object_reader_t code_object_reader, const char *options, hsa_loaded_code_object_t *loaded_code_object) { + TRY IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -1979,6 +2218,7 @@ hsa_status_t hsa_executable_load_program_code_object( {reinterpret_cast(wrapper->code_object_memory)}; return exec->LoadCodeObject( {0}, code_object, options, loaded_code_object); + CATCH } hsa_status_t hsa_executable_load_agent_code_object( @@ -1987,6 +2227,7 @@ hsa_status_t hsa_executable_load_agent_code_object( hsa_code_object_reader_t code_object_reader, const char *options, hsa_loaded_code_object_t *loaded_code_object) { + TRY IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -2004,11 +2245,13 @@ hsa_status_t hsa_executable_load_agent_code_object( {reinterpret_cast(wrapper->code_object_memory)}; return exec->LoadCodeObject( agent, code_object, options, loaded_code_object); + CATCH } hsa_status_t hsa_executable_freeze( hsa_executable_t executable, const char *options) { + TRY IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -2017,12 +2260,14 @@ hsa_status_t hsa_executable_freeze( } return exec->Freeze(options); + CATCH } hsa_status_t hsa_executable_get_info( hsa_executable_t executable, hsa_executable_info_t attribute, void *value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -2032,12 +2277,14 @@ hsa_status_t hsa_executable_get_info( } return exec->GetInfo(attribute, value); + CATCH } hsa_status_t hsa_executable_global_variable_define( hsa_executable_t executable, const char *variable_name, void *address) { + TRY IS_OPEN(); IS_BAD_PTR(variable_name); @@ -2047,6 +2294,7 @@ hsa_status_t hsa_executable_global_variable_define( } return exec->DefineProgramExternalVariable(variable_name, address); + CATCH } hsa_status_t hsa_executable_agent_global_variable_define( @@ -2054,6 +2302,7 @@ hsa_status_t hsa_executable_agent_global_variable_define( hsa_agent_t agent, const char *variable_name, void *address) { + TRY IS_OPEN(); IS_BAD_PTR(variable_name); @@ -2064,6 +2313,7 @@ hsa_status_t hsa_executable_agent_global_variable_define( return exec->DefineAgentExternalVariable( variable_name, agent, HSA_VARIABLE_SEGMENT_GLOBAL, address); + CATCH } hsa_status_t hsa_executable_readonly_variable_define( @@ -2071,6 +2321,7 @@ hsa_status_t hsa_executable_readonly_variable_define( hsa_agent_t agent, const char *variable_name, void *address) { + TRY IS_OPEN(); IS_BAD_PTR(variable_name); @@ -2081,11 +2332,13 @@ hsa_status_t hsa_executable_readonly_variable_define( return exec->DefineAgentExternalVariable( variable_name, agent, HSA_VARIABLE_SEGMENT_READONLY, address); + CATCH } hsa_status_t hsa_executable_validate( hsa_executable_t executable, uint32_t *result) { + TRY IS_OPEN(); IS_BAD_PTR(result); @@ -2095,25 +2348,30 @@ hsa_status_t hsa_executable_validate( } return exec->Validate(result); + CATCH } hsa_status_t hsa_executable_validate_alt( hsa_executable_t executable, const char *options, uint32_t *result) { + TRY IS_OPEN(); IS_BAD_PTR(result); return HSA::hsa_executable_validate(executable, result); + CATCH } -/* deprecated */ hsa_status_t hsa_executable_get_symbol( +/* deprecated */ +hsa_status_t hsa_executable_get_symbol( hsa_executable_t executable, const char *module_name, const char *symbol_name, hsa_agent_t agent, int32_t call_convention, hsa_executable_symbol_t *symbol) { + TRY IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -2136,6 +2394,7 @@ hsa_status_t hsa_executable_validate_alt( return HSA::hsa_executable_get_symbol_by_name( executable, mangled_name.c_str(), exec->IsProgramSymbol(mangled_name.c_str()) ? nullptr : &agent, symbol); + CATCH } hsa_status_t hsa_executable_get_symbol_by_name( @@ -2145,6 +2404,7 @@ hsa_status_t hsa_executable_get_symbol_by_name( // of the specification, but seems like a better // approach to distinguish program/agent symbols. hsa_executable_symbol_t *symbol) { + TRY IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -2161,12 +2421,14 @@ hsa_status_t hsa_executable_get_symbol_by_name( *symbol = loader::Symbol::Handle(sym); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_executable_symbol_get_info( hsa_executable_symbol_t executable_symbol, hsa_executable_symbol_info_t attribute, void *value) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -2177,14 +2439,17 @@ hsa_status_t hsa_executable_symbol_get_info( return sym->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; + CATCH } -/* deprecated */ hsa_status_t hsa_executable_iterate_symbols( +/* deprecated */ +hsa_status_t hsa_executable_iterate_symbols( hsa_executable_t executable, hsa_status_t (*callback)(hsa_executable_t executable, hsa_executable_symbol_t symbol, void *data), void *data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); @@ -2194,6 +2459,7 @@ hsa_status_t hsa_executable_symbol_get_info( } return exec->IterateSymbols(callback, data); + CATCH } hsa_status_t hsa_executable_iterate_agent_symbols( @@ -2204,6 +2470,7 @@ hsa_status_t hsa_executable_iterate_agent_symbols( hsa_executable_symbol_t symbol, void *data), void *data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); @@ -2217,6 +2484,7 @@ hsa_status_t hsa_executable_iterate_agent_symbols( } return exec->IterateAgentSymbols(agent, callback, data); + CATCH } hsa_status_t hsa_executable_iterate_program_symbols( @@ -2225,6 +2493,7 @@ hsa_status_t hsa_executable_iterate_program_symbols( hsa_executable_symbol_t symbol, void *data), void *data) { + TRY IS_OPEN(); IS_BAD_PTR(callback); @@ -2234,6 +2503,7 @@ hsa_status_t hsa_executable_iterate_program_symbols( } return exec->IterateProgramSymbols(callback, data); + CATCH } //===--- Runtime Notifications --------------------------------------------===// 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 901866eaf0..27f7f537bc 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 @@ -40,6 +40,10 @@ // //////////////////////////////////////////////////////////////////////////////// +#include +#include +#include + #include "hsakmt.h" #include "core/inc/runtime.h" @@ -109,11 +113,37 @@ static __forceinline bool IsValid(T* ptr) { return (ptr == NULL) ? NULL : ptr->IsValid(); } +#define TRY try { +#define CATCH } catch(...) { return AMD::handleException(); } +#define CATCHRET(RETURN_TYPE) } catch(...) { return AMD::handleExceptionT(); } + namespace AMD { -hsa_status_t - hsa_amd_coherency_get_type(hsa_agent_t agent_handle, - hsa_amd_coherency_type_t* type) { +hsa_status_t handleException() { + try { + throw; + } catch (const std::bad_alloc& e) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } catch (const std::bad_cast& e) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } catch (std::nested_exception& e) { // Rethrow exceptions from callbacks after unwinding HSA. + e.rethrow_nested(); + return HSA_STATUS_ERROR; + } catch (...) { + assert(false && "Unhandled exception."); + abort(); + return HSA_STATUS_ERROR; + } +} + +template static __forceinline T handleExceptionT() { + handleException(); + abort(); + return T(); +} + +hsa_status_t hsa_amd_coherency_get_type(hsa_agent_t agent_handle, hsa_amd_coherency_type_t* type) { + TRY IS_OPEN(); const core::Agent* agent = core::Agent::Convert(agent_handle); @@ -132,10 +162,12 @@ hsa_status_t *type = gpu_agent->current_coherency_type(); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle, - hsa_amd_coherency_type_t type) { + hsa_amd_coherency_type_t type) { + TRY IS_OPEN(); core::Agent* agent = core::Agent::Convert(agent_handle); @@ -158,10 +190,11 @@ hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle, } return HSA_STATUS_SUCCESS; + CATCH } -hsa_status_t - hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) { +hsa_status_t hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) { + TRY IS_OPEN(); if (ptr == NULL) { @@ -173,17 +206,15 @@ hsa_status_t } return core::Runtime::runtime_singleton_->FillMemory(ptr, value, count); + CATCH } -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, - const hsa_signal_t* dep_signals, - hsa_signal_t completion_signal) { - if (dst == NULL || src == NULL) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } +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, const hsa_signal_t* dep_signals, + hsa_signal_t completion_signal) { + TRY + if (dst == NULL || src == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } if ((num_dep_signals == 0 && dep_signals != NULL) || (num_dep_signals > 0 && dep_signals == NULL)) { @@ -215,10 +246,11 @@ hsa_status_t } return HSA_STATUS_SUCCESS; + CATCH } -hsa_status_t - hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable) { +hsa_status_t hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable) { + TRY IS_OPEN(); core::Queue* cmd_queue = core::Queue::Convert(queue); @@ -229,9 +261,11 @@ hsa_status_t AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, (enable != 0)); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_amd_profiling_async_copy_enable(bool enable) { + TRY IS_OPEN(); return core::Runtime::runtime_singleton_->IterateAgent( @@ -240,11 +274,13 @@ hsa_status_t hsa_amd_profiling_async_copy_enable(bool enable) { return core::Agent::Convert(agent_handle)->profiling_enabled(enable); }, reinterpret_cast(&enable)); + CATCH } 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) { + TRY IS_OPEN(); IS_BAD_PTR(time); @@ -267,10 +303,12 @@ hsa_status_t hsa_amd_profiling_get_dispatch_time( gpu_agent->TranslateTime(signal, *time); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_amd_profiling_get_async_copy_time( hsa_signal_t hsa_signal, hsa_amd_profiling_async_copy_time_t* time) { + TRY IS_OPEN(); IS_BAD_PTR(time); @@ -295,12 +333,13 @@ hsa_status_t hsa_amd_profiling_get_async_copy_time( time->start = signal->signal_.start_ts; time->end = signal->signal_.end_ts; return HSA_STATUS_SUCCESS; + CATCH } -hsa_status_t - hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent_handle, - uint64_t agent_tick, - uint64_t* system_tick) { +hsa_status_t hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent_handle, + uint64_t agent_tick, + uint64_t* system_tick) { + TRY IS_OPEN(); IS_BAD_PTR(system_tick); @@ -318,31 +357,28 @@ hsa_status_t *system_tick = gpu_agent->TranslateTime(agent_tick); return HSA_STATUS_SUCCESS; + CATCH } -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, - hsa_wait_state_t wait_hint, - hsa_signal_value_t* satisfying_value) { +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, hsa_wait_state_t wait_hint, + hsa_signal_value_t* satisfying_value) { + TRY // Do not check for signal invalidation. Invalidation may occur during async // signal handler loop and is not an error. for (uint i = 0; i < signal_count; i++) - assert(hsa_signals[i].handle != 0 && - static_cast*>( - core::Signal::Convert(hsa_signals[i]))->IsValid() && - "Invalid signal."); + assert(core::SharedSignal::Convert(hsa_signals[i])->IsValid() && "Invalid signal."); return core::Signal::WaitAny(signal_count, hsa_signals, conds, values, timeout_hint, wait_hint, satisfying_value); + CATCHRET(uint32_t) } -hsa_status_t - hsa_amd_signal_async_handler(hsa_signal_t hsa_signal, - hsa_signal_condition_t cond, - hsa_signal_value_t value, - hsa_amd_signal_handler handler, void* arg) { +hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t hsa_signal, hsa_signal_condition_t cond, + hsa_signal_value_t value, hsa_amd_signal_handler handler, + void* arg) { + TRY IS_OPEN(); core::Signal* signal = core::Signal::Convert(hsa_signal); @@ -352,10 +388,11 @@ hsa_status_t return HSA_STATUS_ERROR_INVALID_SIGNAL; return core::Runtime::runtime_singleton_->SetAsyncSignalHandler( hsa_signal, cond, value, handler, arg); + CATCH } -hsa_status_t - hsa_amd_async_function(void (*callback)(void* arg), void* arg) { +hsa_status_t hsa_amd_async_function(void (*callback)(void* arg), void* arg) { + TRY IS_OPEN(); IS_BAD_PTR(callback); @@ -363,25 +400,28 @@ hsa_status_t return core::Runtime::runtime_singleton_->SetAsyncSignalHandler( null_signal, HSA_SIGNAL_CONDITION_EQ, 0, (hsa_amd_signal_handler)callback, arg); + CATCH } 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) { + TRY IS_OPEN(); IS_BAD_PTR(cu_mask); core::Queue* cmd_queue = core::Queue::Convert(queue); IS_VALID(cmd_queue); return cmd_queue->SetCUMasking(num_cu_mask_count, cu_mask); + CATCH } 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; - + hsa_agent_t* agents, int num_agent, + void** agent_ptr) { + TRY IS_OPEN(); + *agent_ptr = NULL; if (size == 0 || host_ptr == NULL || agent_ptr == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; @@ -397,9 +437,11 @@ hsa_status_t hsa_amd_memory_lock(void* host_ptr, size_t size, core::Runtime::runtime_singleton_->system_regions_fine()[0]); return system_region->Lock(num_agent, agents, host_ptr, size, agent_ptr); + CATCH } hsa_status_t hsa_amd_memory_unlock(void* host_ptr) { + TRY IS_OPEN(); const amd::MemoryRegion* system_region = @@ -407,12 +449,12 @@ hsa_status_t hsa_amd_memory_unlock(void* host_ptr) { core::Runtime::runtime_singleton_->system_regions_fine()[0]); return system_region->Unlock(host_ptr); + CATCH } -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) { +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) { + TRY IS_OPEN(); IS_BAD_PTR(value); @@ -423,12 +465,14 @@ hsa_status_t } return mem_region->GetPoolInfo(attribute, value); + CATCH } 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) { + TRY IS_OPEN(); IS_BAD_PTR(callback); const core::Agent* agent = core::Agent::Convert(agent_handle); @@ -446,11 +490,12 @@ hsa_status_t hsa_amd_agent_iterate_memory_pools( reinterpret_cast( callback), data); + CATCH } -hsa_status_t - hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size, - uint32_t flags, void** ptr) { +hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size, + uint32_t flags, void** ptr) { + TRY IS_OPEN(); if (size == 0 || ptr == NULL) { @@ -466,15 +511,16 @@ hsa_status_t return core::Runtime::runtime_singleton_->AllocateMemory( mem_region, size, core::MemoryRegion::AllocateRestrict, ptr); + CATCH } hsa_status_t hsa_amd_memory_pool_free(void* ptr) { return HSA::hsa_memory_free(ptr); } -hsa_status_t - hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents, - const uint32_t* flags, const void* ptr) { +hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents, + const uint32_t* flags, const void* ptr) { + TRY IS_OPEN(); if (num_agents == 0 || agents == NULL || flags != NULL || ptr == NULL) { @@ -483,12 +529,12 @@ hsa_status_t return core::Runtime::runtime_singleton_->AllowAccess(num_agents, agents, ptr); + CATCH } -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) { +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) { + TRY IS_OPEN(); if (result == NULL) { @@ -512,11 +558,13 @@ hsa_status_t } return src_mem_region->CanMigrate(*dst_mem_region, *result); + CATCH } hsa_status_t hsa_amd_memory_migrate(const void* ptr, - hsa_amd_memory_pool_t memory_pool, - uint32_t flags) { + hsa_amd_memory_pool_t memory_pool, + uint32_t flags) { + TRY IS_OPEN(); if (ptr == NULL || flags != 0) { @@ -532,11 +580,13 @@ hsa_status_t hsa_amd_memory_migrate(const void* ptr, } return dst_mem_region->Migrate(flags, ptr); + CATCH } 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) { + TRY IS_OPEN(); if (value == NULL) { @@ -555,6 +605,7 @@ hsa_status_t hsa_amd_agent_memory_pool_get_info( } return mem_region->GetAgentPoolInfo(*agent, attribute, value); + CATCH } hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents, @@ -563,6 +614,7 @@ hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents, void** ptr, size_t* metadata_size, const void** metadata) { static const int tinyArraySize=8; + TRY IS_OPEN(); IS_BAD_PTR(agents); IS_BAD_PTR(size); @@ -589,39 +641,49 @@ hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents, if (num_agents > tinyArraySize) delete[] core_agents; return ret; + CATCH } hsa_status_t hsa_amd_interop_unmap_buffer(void* ptr) { + TRY IS_OPEN(); if (ptr != NULL) core::Runtime::runtime_singleton_->InteropUnmap(ptr); return HSA_STATUS_SUCCESS; + CATCH } hsa_status_t hsa_amd_pointer_info(void* ptr, hsa_amd_pointer_info_t* info, void* (*alloc)(size_t), - uint32_t* num_accessible, hsa_agent_t** accessible) { + uint32_t* num_accessible, hsa_agent_t** accessible) { + TRY IS_OPEN(); IS_BAD_PTR(ptr); IS_BAD_PTR(info); return core::Runtime::runtime_singleton_->PtrInfo(ptr, info, alloc, num_accessible, accessible); + CATCH } hsa_status_t hsa_amd_pointer_info_set_userdata(void* ptr, void* userdata) { + TRY IS_OPEN(); IS_BAD_PTR(ptr); return core::Runtime::runtime_singleton_->SetPtrInfoData(ptr, userdata); + CATCH } hsa_status_t hsa_amd_ipc_memory_create(void* ptr, size_t len, hsa_amd_ipc_memory_t* handle) { + TRY IS_OPEN(); IS_BAD_PTR(ptr); IS_BAD_PTR(handle); return core::Runtime::runtime_singleton_->IPCCreate(ptr, len, handle); + CATCH } hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* ipc, size_t len, uint32_t num_agents, const hsa_agent_t* mapping_agents, void** mapped_ptr) { static const int tinyArraySize = 8; + TRY IS_OPEN(); IS_BAD_PTR(mapped_ptr); if (num_agents != 0) IS_BAD_PTR(mapping_agents); @@ -644,12 +706,15 @@ hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* ipc, size_t l return core::Runtime::runtime_singleton_->IPCAttach(ipc, len, num_agents, core_agents, mapped_ptr); + CATCH } hsa_status_t hsa_amd_ipc_memory_detach(void* mapped_ptr) { + TRY IS_OPEN(); IS_BAD_PTR(mapped_ptr); return core::Runtime::runtime_singleton_->IPCDetach(mapped_ptr); + CATCH } } // end of AMD namespace diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp index 84d20cf28d..7f80164996 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp @@ -68,9 +68,8 @@ int InterruptSignal::rtti_id_ = 0; void InterruptSignal::DestroyEvent(HsaEvent* evt) { hsaKmtDestroyEvent(evt); } -InterruptSignal::InterruptSignal(hsa_signal_value_t initial_value, - HsaEvent* use_event) - : Signal(initial_value) { +InterruptSignal::InterruptSignal(hsa_signal_value_t initial_value, HsaEvent* use_event) + : LocalSignal(initial_value), Signal(signal()) { if (use_event != NULL) { event_ = use_event; free_event_ = false;