diff --git a/runtime/hsa-runtime/CMakeLists.txt b/runtime/hsa-runtime/CMakeLists.txt index d9a98c972d..a1cdfc716d 100644 --- a/runtime/hsa-runtime/CMakeLists.txt +++ b/runtime/hsa-runtime/CMakeLists.txt @@ -133,6 +133,7 @@ set ( SRCS "core/util/lnx/os_linux.cpp" "core/runtime/hsa_ext_amd.cpp" "core/runtime/hsa_ext_interface.cpp" "core/runtime/interrupt_signal.cpp" + "core/runtime/ipc_signal.cpp" "core/runtime/isa.cpp" "core/runtime/runtime.cpp" "core/runtime/signal.cpp" diff --git a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp b/runtime/hsa-runtime/core/common/hsa_table_interface.cpp index 2dc5cc8d56..e593f2d4b0 100644 --- a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp +++ b/runtime/hsa-runtime/core/common/hsa_table_interface.cpp @@ -1075,3 +1075,22 @@ hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* ipc, size_t l hsa_status_t hsa_amd_ipc_memory_detach(void* mapped_ptr) { return amdExtTable->hsa_amd_ipc_memory_detach_fn(mapped_ptr); } + +// Mirrors Amd Extension Apis +hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, + const hsa_agent_t* consumers, uint64_t attributes, + hsa_signal_t* signal) { + return amdExtTable->hsa_amd_signal_create_fn(initial_value, num_consumers, consumers, attributes, + signal); +} + +// Mirrors Amd Extension Apis +hsa_status_t HSA_API hsa_amd_ipc_signal_create(hsa_signal_t signal, hsa_amd_ipc_signal_t* handle) { + return amdExtTable->hsa_amd_ipc_signal_create_fn(signal, handle); +} + +// Mirrors Amd Extension Apis +hsa_status_t HSA_API hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle, + hsa_signal_t* signal) { + return amdExtTable->hsa_amd_ipc_signal_attach_fn(handle, signal); +} diff --git a/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/runtime/hsa-runtime/core/inc/amd_aql_queue.h index 08e65866c8..0213a02b4b 100644 --- a/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -339,6 +339,9 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig protected: bool _IsA(rtti_t id) const override { return id == &rtti_id_; } + /// @brief Disallow destroying doorbell apart from its queue. + void doDestroySignal() override { assert(false); } + private: uint32_t ComputeRingBufferMinPkts(); uint32_t ComputeRingBufferMaxPkts(); diff --git a/runtime/hsa-runtime/core/inc/default_signal.h b/runtime/hsa-runtime/core/inc/default_signal.h index 6f3377b28c..20f0f5bb21 100644 --- a/runtime/hsa-runtime/core/inc/default_signal.h +++ b/runtime/hsa-runtime/core/inc/default_signal.h @@ -51,21 +51,18 @@ namespace core { -/// @brief Simple pure memory based signal. +/// @brief Operations for a simple pure memory based signal. /// @brief See base class Signal. -class DefaultSignal : private LocalSignal, public Signal { +class BusyWaitSignal : public Signal { public: - /// @brief Determines if a Signal* can be safely converted to DefaultSignal* + /// @brief Determines if a Signal* can be safely converted to BusyWaitSignal* /// via static_cast. static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id_); } /// @brief See base class Signal. - explicit DefaultSignal(hsa_signal_value_t initial_value); - - /// @brief See base class Signal. - ~DefaultSignal(); + explicit BusyWaitSignal(SharedSignal* abi_block, bool enableIPC); // Below are various methods corresponding to the APIs, which load/store the // signal value or modify the existing signal value automically and with @@ -158,6 +155,29 @@ class DefaultSignal : private LocalSignal, public Signal { protected: bool _IsA(rtti_t id) const { return id == &rtti_id_; } + private: + static int rtti_id_; + + DISALLOW_COPY_AND_ASSIGN(BusyWaitSignal); +}; + +/// @brief Simple memory only signal using a new ABI block. +class DefaultSignal : private LocalSignal, public BusyWaitSignal { + public: + /// @brief Determines if a Signal* can be safely converted to BusyWaitSignal* + /// via static_cast. + static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id_); } + + /// @brief See base class Signal. + explicit DefaultSignal(hsa_signal_value_t initial_value, bool enableIPC = false) + : LocalSignal(initial_value), BusyWaitSignal(signal(), enableIPC) {} + + protected: + bool _IsA(rtti_t id) const { + if (id == &rtti_id_) return true; + return BusyWaitSignal::_IsA(id); + } + private: static int rtti_id_; diff --git a/runtime/hsa-runtime/core/inc/exceptions.h b/runtime/hsa-runtime/core/inc/exceptions.h new file mode 100644 index 0000000000..4e4c97bbcf --- /dev/null +++ b/runtime/hsa-runtime/core/inc/exceptions.h @@ -0,0 +1,67 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// The University of Illinois/NCSA +// Open Source License (NCSA) +// +// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +// +// Developed by: +// +// AMD Research and AMD HSA Software Development +// +// Advanced Micro Devices, Inc. +// +// www.amd.com +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to +// deal with the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimers. +// - Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimers in +// the documentation and/or other materials provided with the distribution. +// - Neither the names of Advanced Micro Devices, Inc, +// nor the names of its contributors may be used to endorse or promote +// products derived from this Software without specific prior written +// permission. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +// DEALINGS WITH THE SOFTWARE. +// +//////////////////////////////////////////////////////////////////////////////// + +#ifndef HSA_RUNTIME_CORE_INC_EXCEPTIONS_H +#define HSA_RUNTIME_CORE_INC_EXCEPTIONS_H + +#include +#include + +#include "core/inc/hsa_internal.h" + +namespace AMD { + +/// @brief Exception type which carries an error code to return to the user. +class hsa_exception : public std::exception { + public: + hsa_exception(hsa_status_t error, const char* description) : err_(error), desc_(description) {} + hsa_status_t error_code() const noexcept { return err_; } + const char* what() const noexcept override { return desc_.c_str(); } + + private: + hsa_status_t err_; + std::string desc_; +}; + +} // namespace AMD + +#endif // HSA_RUNTIME_CORE_INC_EXCEPTIONS_H diff --git a/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h b/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h index 763e8553b1..94024a9206 100644 --- a/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h +++ b/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h @@ -94,6 +94,11 @@ hsa_status_t HSA_API hsa_status_t HSA_API hsa_amd_async_function(void (*callback)(void* arg), void* arg); +// Mirrors Amd Extension Apis +hsa_status_t HSA_API hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, + const hsa_agent_t* consumers, uint64_t attributes, + hsa_signal_t* signal); + // Mirrors Amd Extension Apis uint32_t HSA_API hsa_amd_signal_wait_any(uint32_t signal_count, hsa_signal_t* signals, @@ -199,6 +204,12 @@ hsa_status_t hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* handle, size_ // Mirrors Amd Extension Apis hsa_status_t hsa_amd_ipc_memory_detach(void* mapped_ptr); +// Mirrors Amd Extension Apis +hsa_status_t hsa_amd_ipc_signal_create(hsa_signal_t signal, hsa_amd_ipc_signal_t* handle); + +// Mirrors Amd Extension Apis +hsa_status_t hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle, hsa_signal_t* signal); + } // end of AMD namespace #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/ipc_signal.h b/runtime/hsa-runtime/core/inc/ipc_signal.h new file mode 100644 index 0000000000..9303b79e55 --- /dev/null +++ b/runtime/hsa-runtime/core/inc/ipc_signal.h @@ -0,0 +1,112 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// The University of Illinois/NCSA +// Open Source License (NCSA) +// +// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +// +// Developed by: +// +// AMD Research and AMD HSA Software Development +// +// Advanced Micro Devices, Inc. +// +// www.amd.com +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to +// deal with the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimers. +// - Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimers in +// the documentation and/or other materials provided with the distribution. +// - Neither the names of Advanced Micro Devices, Inc, +// nor the names of its contributors may be used to endorse or promote +// products derived from this Software without specific prior written +// permission. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +// DEALINGS WITH THE SOFTWARE. +// +//////////////////////////////////////////////////////////////////////////////// + +#ifndef HSA_RUNTME_CORE_INC_IPC_SIGNAL_H_ +#define HSA_RUNTME_CORE_INC_IPC_SIGNAL_H_ + +#include +#include + +#include "core/inc/signal.h" +#include "core/inc/default_signal.h" +#include "core/util/locks.h" + +namespace core { + +/// @brief Container for ipc shared memory. +class SharedMemory { + public: + SharedMemory(const hsa_amd_ipc_memory_t* handle, size_t len); + ~SharedMemory(); + SharedMemory(SharedMemory&&); + + void* ptr() const { return ptr_; } + + private: + void* ptr_; +}; + +/// @brief Container for ipc signal abi block. +class SharedMemorySignal { + public: + explicit SharedMemorySignal(const hsa_amd_ipc_memory_t* handle) : signal_(handle, 4096) { + if (!signal()->IsValid()) + throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT, "IPC Signal handle is invalid."); + } + SharedSignal* signal() const { return reinterpret_cast(signal_.ptr()); } + + private: + SharedMemory signal_; +}; + +/// @brief Memory only signal using a shared memory ABI block. +class IPCSignal : private SharedMemorySignal, public BusyWaitSignal { + public: + /// @brief Creates a sharable handle for an IPC enabled signal. + static void CreateHandle(Signal* signal, hsa_amd_ipc_signal_t* ipc_handle); + + /// @brief Opens an IPC signal from its IPC handle. + static Signal* Attach(const hsa_amd_ipc_signal_t* ipc_handle); + + /// @brief Determines if a Signal* can be safely converted to BusyWaitSignal* + /// via static_cast. + static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id_); } + + protected: + bool _IsA(rtti_t id) const { + if (id == &rtti_id_) return true; + return BusyWaitSignal::_IsA(id); + } + + private: + static int rtti_id_; + static KernelMutex lock_; + + explicit IPCSignal(SharedMemorySignal&& abi_block) + : SharedMemorySignal(std::move(abi_block)), BusyWaitSignal(signal(), true) {} + + DISALLOW_COPY_AND_ASSIGN(IPCSignal); +}; + +} // namespace core + +#endif // HSA_RUNTME_CORE_INC_IPC_SIGNAL_H_ diff --git a/runtime/hsa-runtime/core/inc/signal.h b/runtime/hsa-runtime/core/inc/signal.h index 502a7880c0..028138a9a1 100644 --- a/runtime/hsa-runtime/core/inc/signal.h +++ b/runtime/hsa-runtime/core/inc/signal.h @@ -45,12 +45,15 @@ #ifndef HSA_RUNTME_CORE_INC_SIGNAL_H_ #define HSA_RUNTME_CORE_INC_SIGNAL_H_ +#include + #include "hsakmt.h" #include "core/common/shared.h" #include "core/inc/runtime.h" #include "core/inc/checked.h" +#include "core/inc/exceptions.h" #include "core/util/utils.h" #include "core/util/locks.h" @@ -67,20 +70,28 @@ struct SharedSignal { Signal* core_signal; Check<0x71FCCA6A3D5D5276, true> id; - bool IsValid() const { return id.IsValid(); } - SharedSignal() { + memset(&amd_signal, 0, sizeof(amd_signal)); amd_signal.kind = AMD_SIGNAL_KIND_INVALID; core_signal = nullptr; } + bool IsValid() const { return (Convert(this).handle != 0) && id.IsValid(); } + + bool IsIPC() const { return 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 __forceinline hsa_signal_t Convert(const SharedSignal* signal) { + assert(signal != nullptr && "Conversion on null Signal object."); + const uint64_t handle = static_cast(reinterpret_cast(&signal->amd_signal)); + const hsa_signal_t signal_handle = {handle}; + return signal_handle; + } }; static_assert(std::is_standard_layout::value, "SharedSignal must remain standard layout for IPC use."); @@ -107,49 +118,86 @@ class LocalSignal { class Signal { public: /// @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(); + explicit Signal(SharedSignal* abi_block, bool enableIPC = false) + : signal_(abi_block->amd_signal), async_copy_agent_(NULL), refcount_(1) { + assert(abi_block != nullptr && "Signal abi_block must not be NULL"); - invalid_ = false; waiting_ = 0; - retained_ = 0; + retained_ = 1; - abi_block->core_signal = this; + if (enableIPC) { + abi_block->core_signal = nullptr; + registerIpc(); + } else { + abi_block->core_signal = this; + } } - virtual ~Signal() { invalid_ = true; } + /// @brief Interface to discard a signal handle (hsa_signal_t) + /// Decrements signal ref count and invokes doDestroySignal() when + /// Signal is no longer in use. + void DestroySignal() { + // If handle is now invalid wake any retained sleepers. + if (--refcount_ == 0) CasRelaxed(0, 0); + // Release signal, last release will destroy the object. + Release(); + } - bool IsValid() const { return !invalid_; } - - /// @brief Converts from this implementation class to the public + /// @brief Converts from this interface class to the public /// hsa_signal_t type - an opaque handle. static __forceinline hsa_signal_t Convert(Signal* signal) { - assert(signal != nullptr && signal->IsValid() && "Conversion on invalid Signal object."); + assert(signal != nullptr && "Conversion on null Signal object."); const uint64_t handle = static_cast(reinterpret_cast(&signal->signal_)); const hsa_signal_t signal_handle = {handle}; return signal_handle; } - /// @brief Converts from this implementation class to the public + /// @brief Converts from this interface class to the public /// hsa_signal_t type - an opaque handle. static __forceinline const hsa_signal_t Convert(const Signal* signal) { - assert(signal != nullptr && signal->IsValid() && "Conversion on invalid Signal object."); + assert(signal != nullptr && "Conversion on null Signal object."); const uint64_t handle = static_cast(reinterpret_cast(&signal->signal_)); const hsa_signal_t signal_handle = {handle}; return signal_handle; } /// @brief Converts from public hsa_signal_t type (an opaque handle) to - /// this implementation class object. + /// this interface class object. static __forceinline Signal* Convert(hsa_signal_t signal) { + if (signal.handle == 0) + throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT, "Signal handle is invalid."); SharedSignal* shared = SharedSignal::Convert(signal); - if (shared->core_signal != nullptr) + if (!shared->IsValid()) + throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_SIGNAL, "Signal handle is invalid."); + if (shared->IsIPC()) { + Signal* ret = lookupIpc(signal); + if (ret == nullptr) + throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_SIGNAL, "Signal handle is invalid."); + return ret; + } else { return shared->core_signal; - else - throw std::bad_cast(); + } } + static Signal* DuplicateHandle(hsa_signal_t signal) { + if (signal.handle == 0) return nullptr; + SharedSignal* shared = SharedSignal::Convert(signal); + + if (!shared->IsIPC()) { + if (!shared->IsValid()) return nullptr; + shared->core_signal->refcount_++; + shared->core_signal->Retain(); + return shared->core_signal; + } + + // IPC signals may only be duplicated while holding the ipcMap lock. + return duplicateIpc(signal); + } + + bool IsValid() const { return refcount_ != 0; } + + bool __forceinline isIPC() const { return SharedSignal::Convert(Convert(this))->IsIPC(); } + // Below are various methods corresponding to the APIs, which load/store the // signal value or modify the existing signal value automically and with // specified memory ordering semantics. @@ -230,13 +278,9 @@ class Signal { __forceinline bool IsType(rtti_t id) { return _IsA(id); } - /// @brief Allows special case interaction with signal destruction cleanup. - void Retain() { atomic::Increment(&retained_); } - void Release() { atomic::Decrement(&retained_); } - - /// @brief Checks if signal is currently in use such that it should not be - /// deleted. - bool InUse() const { return (retained_ != 0) || (waiting_ != 0); } + /// @brief Prevents the signal from being destroyed until the matching Release(). + void Retain() { retained_++; } + void Release(); /// @brief Checks if signal is currently in use by a wait API. bool InWaiting() const { return waiting_ != 0; } @@ -253,25 +297,40 @@ class Signal { amd_signal_t& signal_; protected: + virtual ~Signal(); + + /// @brief Overrideable deletion function + virtual void doDestroySignal() { delete this; } + /// @brief Simple RTTI type checking helper /// Returns true if the object can be converted to the query type via /// static_cast. /// Do not use directly. Use IsType in the desired derived type instead. virtual bool _IsA(rtti_t id) const = 0; - /// @variable Indicates if signal is valid or not. - volatile bool invalid_; - /// @variable Indicates number of runtime threads waiting on this signal. /// Value of zero means no waits. - volatile uint32_t waiting_; - - volatile uint32_t retained_; + std::atomic waiting_; /// @variable Pointer to agent used to perform an async copy. core::Agent* async_copy_agent_; private: + static KernelMutex ipcLock_; + static std::map ipcMap_; + + static Signal* lookupIpc(hsa_signal_t signal); + static Signal* duplicateIpc(hsa_signal_t signal); + + /// @variable Ref count of this signal's handle (see IPC APIs) + std::atomic refcount_; + + /// @variable Count of handle references and Retain() calls for this handle (see IPC APIs) + std::atomic retained_; + + void registerIpc(); + bool deregisterIpc(); + DISALLOW_COPY_AND_ASSIGN(Signal); }; diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index d11524954e..0298266a5a 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -582,6 +582,7 @@ void GpuAgent::InitDma() { if (blits_[BlitDevToHost] == NULL) { // Share utility queue with device-to-host blits. + if (queues_[QueueUtility] == nullptr) queues_[QueueUtility] = CreateInterceptibleQueue(); blits_[BlitDevToHost] = CreateBlitKernel(queues_[QueueUtility]); assert(blits_[BlitDevToHost] != NULL && "Blit creation failed"); } @@ -597,14 +598,15 @@ hsa_status_t GpuAgent::PostToolsInit() { BindTrapHandler(); // Defer utility queue creation to allow tools to intercept. - queues_[QueueUtility] = CreateInterceptibleQueue(); + if (queues_[QueueUtility] == nullptr) queues_[QueueUtility] = CreateInterceptibleQueue(); if (queues_[QueueUtility] == NULL) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } // Share utility queue with device-to-device blits. - blits_[BlitDevToDev] = CreateBlitKernel(queues_[QueueUtility]); + if (blits_[BlitDevToDev] == nullptr) + blits_[BlitDevToDev] = CreateBlitKernel(queues_[QueueUtility]); if (blits_[BlitDevToDev] == NULL) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; diff --git a/runtime/hsa-runtime/core/runtime/default_signal.cpp b/runtime/hsa-runtime/core/runtime/default_signal.cpp index b1a89a5af7..249736dd79 100644 --- a/runtime/hsa-runtime/core/runtime/default_signal.cpp +++ b/runtime/hsa-runtime/core/runtime/default_signal.cpp @@ -46,49 +46,45 @@ namespace core { int DefaultSignal::rtti_id_ = 0; +int BusyWaitSignal::rtti_id_ = 0; -DefaultSignal::DefaultSignal(hsa_signal_value_t initial_value) - : LocalSignal(initial_value), Signal(signal()) { +BusyWaitSignal::BusyWaitSignal(SharedSignal* abi_block, bool enableIPC) + : Signal(abi_block, enableIPC) { signal_.kind = AMD_SIGNAL_KIND_USER; signal_.event_mailbox_ptr = NULL; - HSA::hsa_memory_register(this, sizeof(DefaultSignal)); } -DefaultSignal::~DefaultSignal() { - invalid_ = true; - while (InUse()) - ; - HSA::hsa_memory_deregister(this, sizeof(DefaultSignal)); -} - -hsa_signal_value_t DefaultSignal::LoadRelaxed() { +hsa_signal_value_t BusyWaitSignal::LoadRelaxed() { return hsa_signal_value_t( atomic::Load(&signal_.value, std::memory_order_relaxed)); } -hsa_signal_value_t DefaultSignal::LoadAcquire() { +hsa_signal_value_t BusyWaitSignal::LoadAcquire() { return hsa_signal_value_t( atomic::Load(&signal_.value, std::memory_order_acquire)); } -void DefaultSignal::StoreRelaxed(hsa_signal_value_t value) { +void BusyWaitSignal::StoreRelaxed(hsa_signal_value_t value) { atomic::Store(&signal_.value, int64_t(value), std::memory_order_relaxed); } -void DefaultSignal::StoreRelease(hsa_signal_value_t value) { +void BusyWaitSignal::StoreRelease(hsa_signal_value_t value) { atomic::Store(&signal_.value, int64_t(value), std::memory_order_release); } -hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition, - hsa_signal_value_t compare_value, - uint64_t timeout, - hsa_wait_state_t wait_hint) { - atomic::Increment(&waiting_); - MAKE_SCOPE_GUARD([&]() { atomic::Decrement(&waiting_); }); +hsa_signal_value_t BusyWaitSignal::WaitRelaxed(hsa_signal_condition_t condition, + hsa_signal_value_t compare_value, uint64_t timeout, + hsa_wait_state_t wait_hint) { + Retain(); + MAKE_SCOPE_GUARD([&]() { Release(); }); + + waiting_++; + MAKE_SCOPE_GUARD([&]() { waiting_--; }); bool condition_met = false; int64_t value; - debug_warning(!g_use_interrupt_wait && "Use of non-host signal in host signal wait API."); + debug_warning((!g_use_interrupt_wait || isIPC()) && + "Use of non-host signal in host signal wait API."); timer::fast_clock::time_point start_time, time; start_time = timer::fast_clock::now(); @@ -100,7 +96,7 @@ hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition, double(timeout) / double(hsa_freq)); while (true) { - if (invalid_) return 0; + if (!IsValid()) return 0; value = atomic::Load(&signal_.value, std::memory_order_relaxed); @@ -135,139 +131,138 @@ hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition, } } -hsa_signal_value_t DefaultSignal::WaitAcquire(hsa_signal_condition_t condition, - hsa_signal_value_t compare_value, - uint64_t timeout, - hsa_wait_state_t wait_hint) { +hsa_signal_value_t BusyWaitSignal::WaitAcquire(hsa_signal_condition_t condition, + hsa_signal_value_t compare_value, uint64_t timeout, + hsa_wait_state_t wait_hint) { hsa_signal_value_t ret = WaitRelaxed(condition, compare_value, timeout, wait_hint); std::atomic_thread_fence(std::memory_order_acquire); return ret; } -void DefaultSignal::AndRelaxed(hsa_signal_value_t value) { +void BusyWaitSignal::AndRelaxed(hsa_signal_value_t value) { atomic::And(&signal_.value, int64_t(value), std::memory_order_relaxed); } -void DefaultSignal::AndAcquire(hsa_signal_value_t value) { +void BusyWaitSignal::AndAcquire(hsa_signal_value_t value) { atomic::And(&signal_.value, int64_t(value), std::memory_order_acquire); } -void DefaultSignal::AndRelease(hsa_signal_value_t value) { +void BusyWaitSignal::AndRelease(hsa_signal_value_t value) { atomic::And(&signal_.value, int64_t(value), std::memory_order_release); } -void DefaultSignal::AndAcqRel(hsa_signal_value_t value) { +void BusyWaitSignal::AndAcqRel(hsa_signal_value_t value) { atomic::And(&signal_.value, int64_t(value), std::memory_order_acq_rel); } -void DefaultSignal::OrRelaxed(hsa_signal_value_t value) { +void BusyWaitSignal::OrRelaxed(hsa_signal_value_t value) { atomic::Or(&signal_.value, int64_t(value), std::memory_order_relaxed); } -void DefaultSignal::OrAcquire(hsa_signal_value_t value) { +void BusyWaitSignal::OrAcquire(hsa_signal_value_t value) { atomic::Or(&signal_.value, int64_t(value), std::memory_order_acquire); } -void DefaultSignal::OrRelease(hsa_signal_value_t value) { +void BusyWaitSignal::OrRelease(hsa_signal_value_t value) { atomic::Or(&signal_.value, int64_t(value), std::memory_order_release); } -void DefaultSignal::OrAcqRel(hsa_signal_value_t value) { +void BusyWaitSignal::OrAcqRel(hsa_signal_value_t value) { atomic::Or(&signal_.value, int64_t(value), std::memory_order_acq_rel); } -void DefaultSignal::XorRelaxed(hsa_signal_value_t value) { +void BusyWaitSignal::XorRelaxed(hsa_signal_value_t value) { atomic::Xor(&signal_.value, int64_t(value), std::memory_order_relaxed); } -void DefaultSignal::XorAcquire(hsa_signal_value_t value) { +void BusyWaitSignal::XorAcquire(hsa_signal_value_t value) { atomic::Xor(&signal_.value, int64_t(value), std::memory_order_acquire); } -void DefaultSignal::XorRelease(hsa_signal_value_t value) { +void BusyWaitSignal::XorRelease(hsa_signal_value_t value) { atomic::Xor(&signal_.value, int64_t(value), std::memory_order_release); } -void DefaultSignal::XorAcqRel(hsa_signal_value_t value) { +void BusyWaitSignal::XorAcqRel(hsa_signal_value_t value) { atomic::Xor(&signal_.value, int64_t(value), std::memory_order_acq_rel); } -void DefaultSignal::AddRelaxed(hsa_signal_value_t value) { +void BusyWaitSignal::AddRelaxed(hsa_signal_value_t value) { atomic::Add(&signal_.value, int64_t(value), std::memory_order_relaxed); } -void DefaultSignal::AddAcquire(hsa_signal_value_t value) { +void BusyWaitSignal::AddAcquire(hsa_signal_value_t value) { atomic::Add(&signal_.value, int64_t(value), std::memory_order_acquire); } -void DefaultSignal::AddRelease(hsa_signal_value_t value) { +void BusyWaitSignal::AddRelease(hsa_signal_value_t value) { atomic::Add(&signal_.value, int64_t(value), std::memory_order_release); } -void DefaultSignal::AddAcqRel(hsa_signal_value_t value) { +void BusyWaitSignal::AddAcqRel(hsa_signal_value_t value) { atomic::Add(&signal_.value, int64_t(value), std::memory_order_acq_rel); } -void DefaultSignal::SubRelaxed(hsa_signal_value_t value) { +void BusyWaitSignal::SubRelaxed(hsa_signal_value_t value) { atomic::Sub(&signal_.value, int64_t(value), std::memory_order_relaxed); } -void DefaultSignal::SubAcquire(hsa_signal_value_t value) { +void BusyWaitSignal::SubAcquire(hsa_signal_value_t value) { atomic::Sub(&signal_.value, int64_t(value), std::memory_order_acquire); } -void DefaultSignal::SubRelease(hsa_signal_value_t value) { +void BusyWaitSignal::SubRelease(hsa_signal_value_t value) { atomic::Sub(&signal_.value, int64_t(value), std::memory_order_release); } -void DefaultSignal::SubAcqRel(hsa_signal_value_t value) { +void BusyWaitSignal::SubAcqRel(hsa_signal_value_t value) { atomic::Sub(&signal_.value, int64_t(value), std::memory_order_acq_rel); } -hsa_signal_value_t DefaultSignal::ExchRelaxed(hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::ExchRelaxed(hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value), std::memory_order_relaxed)); } -hsa_signal_value_t DefaultSignal::ExchAcquire(hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::ExchAcquire(hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value), std::memory_order_acquire)); } -hsa_signal_value_t DefaultSignal::ExchRelease(hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::ExchRelease(hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value), std::memory_order_release)); } -hsa_signal_value_t DefaultSignal::ExchAcqRel(hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::ExchAcqRel(hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Exchange(&signal_.value, int64_t(value), std::memory_order_acq_rel)); } -hsa_signal_value_t DefaultSignal::CasRelaxed(hsa_signal_value_t expected, - hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::CasRelaxed(hsa_signal_value_t expected, + hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value), int64_t(expected), std::memory_order_relaxed)); } -hsa_signal_value_t DefaultSignal::CasAcquire(hsa_signal_value_t expected, - hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::CasAcquire(hsa_signal_value_t expected, + hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value), int64_t(expected), std::memory_order_acquire)); } -hsa_signal_value_t DefaultSignal::CasRelease(hsa_signal_value_t expected, - hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::CasRelease(hsa_signal_value_t expected, + hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value), int64_t(expected), std::memory_order_release)); } -hsa_signal_value_t DefaultSignal::CasAcqRel(hsa_signal_value_t expected, - hsa_signal_value_t value) { +hsa_signal_value_t BusyWaitSignal::CasAcqRel(hsa_signal_value_t expected, + hsa_signal_value_t value) { return hsa_signal_value_t(atomic::Cas(&signal_.value, int64_t(value), int64_t(expected), std::memory_order_acq_rel)); diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index 890f642852..2be263bffd 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -44,7 +44,6 @@ // This file does argument checking and conversion to C++. #include #include -#include #include #include @@ -66,8 +65,6 @@ #include "core/inc/queue.h" #include "core/inc/signal.h" #include "core/inc/cache.h" -#include "core/inc/default_signal.h" -#include "core/inc/interrupt_signal.h" #include "core/inc/amd_loader_context.hpp" #include "inc/hsa_ven_amd_loader.h" #include "inc/hsa_ven_amd_aqlprofile.h" @@ -207,16 +204,16 @@ namespace HSA { // Init/Shutdown routines //---------------------------------------------------------------------------// hsa_status_t hsa_init() { - TRY + TRY; return core::Runtime::runtime_singleton_->Acquire(); - CATCH + CATCH; } hsa_status_t hsa_shut_down() { - TRY + TRY; IS_OPEN(); return core::Runtime::runtime_singleton_->Release(); - CATCH + CATCH; } //---------------------------------------------------------------------------// @@ -224,15 +221,15 @@ hsa_status_t hsa_shut_down() { //---------------------------------------------------------------------------// hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void* value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); return core::Runtime::runtime_singleton_->GetSystemInfo(attribute, value); - CATCH + CATCH; } hsa_status_t hsa_extension_get_name(uint16_t extension, const char** name) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(name); switch (extension) { @@ -262,13 +259,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 + CATCH; } hsa_status_t hsa_system_extension_supported(uint16_t extension, uint16_t version_major, uint16_t version_minor, bool* result) { - TRY + TRY; IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && @@ -296,12 +293,12 @@ hsa_status_t } return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_system_major_extension_supported(uint16_t extension, uint16_t version_major, uint16_t* version_minor, bool* result) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(version_minor); IS_BAD_PTR(result); @@ -332,7 +329,7 @@ hsa_status_t hsa_system_major_extension_supported(uint16_t extension, uint16_t v *result = false; return HSA_STATUS_SUCCESS; - CATCH + CATCH; } static size_t get_extension_table_length(uint16_t extension, uint16_t major, uint16_t minor) { @@ -390,16 +387,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 + TRY; return HSA::hsa_system_get_major_extension_table( extension, version_major, get_extension_table_length(extension, version_major, version_minor), table); - CATCH + CATCH; } hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t version_major, size_t table_length, void* table) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(table); @@ -484,7 +481,7 @@ hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t v } return HSA_STATUS_ERROR; - CATCH + CATCH; } //---------------------------------------------------------------------------// @@ -493,29 +490,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 + TRY; IS_OPEN(); IS_BAD_PTR(callback); return core::Runtime::runtime_singleton_->IterateAgent(callback, data); - CATCH + CATCH; } hsa_status_t hsa_agent_get_info(hsa_agent_t agent_handle, hsa_agent_info_t attribute, void* value) { - TRY + 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 + CATCH; } hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent_handle, hsa_profile_t profile, uint16_t* mask) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(mask); IS_BAD_PROFILE(profile); @@ -524,36 +521,36 @@ hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent_handle, *mask = 0; return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, void* value) { - TRY + TRY; IS_OPEN(); core::Cache* Cache = core::Cache::Convert(cache); IS_VALID(Cache); IS_BAD_PTR(value); return Cache->GetInfo(attribute, value); - CATCH + 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 + 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 + 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 + TRY; IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && @@ -582,13 +579,13 @@ hsa_status_t } return HSA_STATUS_SUCCESS; - CATCH + 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 + TRY; IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && @@ -614,7 +611,7 @@ hsa_status_t hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t } return HSA_STATUS_SUCCESS; - CATCH + CATCH; } /// @brief Api to create a user mode queue. @@ -639,7 +636,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 + TRY; IS_OPEN(); if ((queue == NULL) || (size == 0) || (!IsPowerOfTwo(size)) || @@ -676,14 +673,14 @@ hsa_status_t hsa_queue_create( } return status; - CATCH + 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 + TRY; IS_OPEN(); if ((queue == NULL) || (region.handle == 0) || @@ -710,7 +707,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 + CATCH; } /// @brief Api to destroy a user mode queue @@ -719,14 +716,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 + 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 + CATCH; } /// @brief Api to inactivate a user mode queue @@ -735,14 +732,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 + 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 + CATCH; } /// @brief Api to read the Read Index of Queue using Acquire semantics @@ -751,11 +748,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadReadIndexAcquire(); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to read the Read Index of Queue using Relaxed semantics @@ -764,11 +761,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadReadIndexRelaxed(); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to read the Write Index of Queue using Acquire semantics @@ -777,11 +774,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadWriteIndexAcquire(); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to read the Write Index of Queue using Relaxed semantics @@ -790,11 +787,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadWriteIndexAcquire(); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to store the Read Index of Queue using Relaxed semantics @@ -804,11 +801,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreReadIndexRelaxed(value); - CATCHRET(void) + CATCHRET(void); } /// @brief Api to store the Read Index of Queue using Release semantics @@ -817,11 +814,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreReadIndexRelease(value); - CATCHRET(void) + CATCHRET(void); } /// @brief Api to store the Write Index of Queue using Relaxed semantics @@ -831,11 +828,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreWriteIndexRelaxed(value); - CATCHRET(void) + CATCHRET(void); } /// @brief Api to store the Write Index of Queue using Release semantics @@ -844,11 +841,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreWriteIndexRelease(value); - CATCHRET(void) + CATCHRET(void); } /// @brief Api to compare and swap the Write Index of Queue using Acquire and @@ -863,11 +860,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcqRel(expected, value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to compare and swap the Write Index of Queue using Acquire @@ -882,11 +879,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcquire(expected, value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to compare and swap the Write Index of Queue using Relaxed @@ -902,11 +899,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexRelaxed(expected, value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to compare and swap the Write Index of Queue using Release @@ -921,11 +918,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexRelease(expected, value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Acquire and Release @@ -937,11 +934,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcqRel(value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Acquire Semantics @@ -952,11 +949,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcquire(value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Relaxed Semantics @@ -968,11 +965,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexRelaxed(value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Release Semantics @@ -983,11 +980,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 + TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexRelease(value); - CATCHRET(uint64_t) + CATCHRET(uint64_t); } //----------------------------------------------------------------------------- @@ -996,19 +993,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 + 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 + CATCH; } hsa_status_t hsa_region_get_info(hsa_region_t region, hsa_region_info_t attribute, void* value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -1016,11 +1013,11 @@ hsa_status_t hsa_region_get_info(hsa_region_t region, IS_VALID(mem_region); return mem_region->GetInfo(attribute, value); - CATCH + CATCH; } hsa_status_t hsa_memory_register(void* address, size_t size) { - TRY + TRY; IS_OPEN(); if (size == 0 && address != NULL) { @@ -1028,20 +1025,20 @@ hsa_status_t hsa_memory_register(void* address, size_t size) { } return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_memory_deregister(void* address, size_t size) { - TRY + TRY; IS_OPEN(); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_memory_allocate(hsa_region_t region, size_t size, void** ptr) { - TRY + TRY; IS_OPEN(); if (size == 0 || ptr == NULL) { @@ -1053,11 +1050,11 @@ hsa_status_t return core::Runtime::runtime_singleton_->AllocateMemory( mem_region, size, core::MemoryRegion::AllocateNoFlags, ptr); - CATCH + CATCH; } hsa_status_t hsa_memory_free(void* ptr) { - TRY + TRY; IS_OPEN(); if (ptr == NULL) { @@ -1065,13 +1062,13 @@ hsa_status_t hsa_memory_free(void* ptr) { } return core::Runtime::runtime_singleton_->FreeMemory(ptr); - CATCH + CATCH; } hsa_status_t hsa_memory_assign_agent(void* ptr, hsa_agent_t agent_handle, hsa_access_permission_t access) { - TRY + TRY; IS_OPEN(); if ((ptr == NULL) || (access < HSA_ACCESS_PERMISSION_RO) || @@ -1083,11 +1080,11 @@ hsa_status_t hsa_memory_assign_agent(void* ptr, IS_VALID(agent); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_memory_copy(void* dst, const void* src, size_t size) { - TRY + TRY; IS_OPEN(); if (dst == NULL || src == NULL) { @@ -1099,114 +1096,59 @@ hsa_status_t hsa_memory_copy(void* dst, const void* src, size_t size) { } return core::Runtime::runtime_singleton_->CopyMemory(dst, src, size); - CATCH + CATCH; } //----------------------------------------------------------------------------- // Signals //----------------------------------------------------------------------------- -struct AgentHandleCompare { - bool operator()(const hsa_agent_t& lhs, const hsa_agent_t& rhs) const { - return lhs.handle < rhs.handle; - } -}; - 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); - - core::Signal* ret; - - bool uses_host = false; - - if (num_consumers > 0) { - IS_BAD_PTR(consumers); - - // Check for duplicates in consumers. - std::set consumer_set = - std::set(consumers, - consumers + num_consumers); - if (consumer_set.size() != num_consumers) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - for (const core::Agent* cpu_agent : - core::Runtime::runtime_singleton_->cpu_agents()) { - uses_host |= - (consumer_set.find(cpu_agent->public_handle()) != consumer_set.end()); - } - } else { - uses_host = true; - } - - if (core::g_use_interrupt_wait && uses_host) { - ret = new core::InterruptSignal(initial_value); - } else { - ret = new core::DefaultSignal(initial_value); - } - CHECK_ALLOC(ret); - - *hsa_signal = core::Signal::Convert(ret); - - if (hsa_signal->handle == 0) { - delete ret; - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - - return HSA_STATUS_SUCCESS; - CATCH + return AMD::hsa_amd_signal_create(initial_value, num_consumers, consumers, 0, hsa_signal); } hsa_status_t hsa_signal_destroy(hsa_signal_t hsa_signal) { - TRY + TRY; IS_OPEN(); - - if (hsa_signal.handle == 0) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - core::Signal* signal = core::Signal::Convert(hsa_signal); - IS_VALID(signal); - delete signal; + signal->DestroySignal(); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_signal_value_t hsa_signal_load_relaxed(hsa_signal_t hsa_signal) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->LoadRelaxed(); - CATCHRET(hsa_signal_value_t) + CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t hsa_signal) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->LoadAcquire(); - CATCHRET(hsa_signal_value_t) + CATCHRET(hsa_signal_value_t); } void hsa_signal_store_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->StoreRelaxed(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_store_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->StoreRelease(value); - CATCHRET(void) + CATCHRET(void); } hsa_signal_value_t @@ -1215,12 +1157,12 @@ hsa_signal_value_t hsa_signal_value_t compare_value, uint64_t timeout_hint, hsa_wait_state_t wait_state_hint) { - TRY + 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) + CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t hsa_signal, @@ -1228,18 +1170,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 + 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) + 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 + 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])); @@ -1252,17 +1194,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 + CATCH; } hsa_status_t hsa_signal_group_destroy(hsa_signal_group_t signal_group) { - TRY + TRY; IS_OPEN(); core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); delete group; return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, @@ -1270,7 +1212,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 + TRY; IS_OPEN(); const core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); @@ -1281,7 +1223,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 + CATCH; } hsa_status_t hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, @@ -1289,247 +1231,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 + 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 + CATCH; } void hsa_signal_and_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndRelaxed(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_and_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcquire(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_and_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndRelease(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_and_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcqRel(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_or_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrRelaxed(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_or_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcquire(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_or_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrRelease(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_or_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcqRel(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_xor_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorRelaxed(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_xor_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcquire(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_xor_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorRelease(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_xor_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcqRel(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_add_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->AddRelaxed(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_add_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcquire(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_add_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddRelease(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_add_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcqRel(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_subtract_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubRelaxed(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_subtract_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcquire(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_subtract_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubRelease(value); - CATCHRET(void) + CATCHRET(void); } void hsa_signal_subtract_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcqRel(value); - CATCHRET(void) + CATCHRET(void); } hsa_signal_value_t hsa_signal_exchange_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchRelaxed(value); - CATCHRET(hsa_signal_value_t) + CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_exchange_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcquire(value); - CATCHRET(hsa_signal_value_t) + CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_exchange_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { - TRY + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchRelease(value); - CATCHRET(hsa_signal_value_t) + 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 + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcqRel(value); - CATCHRET(hsa_signal_value_t) + 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 + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasRelaxed(expected, value); - CATCHRET(hsa_signal_value_t) + 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 + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcquire(expected, value); - CATCHRET(hsa_signal_value_t) + 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 + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasRelease(expected, value); - CATCHRET(hsa_signal_value_t) + 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 + TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcqRel(expected, value); - CATCHRET(hsa_signal_value_t) + CATCHRET(hsa_signal_value_t); } //===--- Instruction Set Architecture -------------------------------------===// @@ -1541,7 +1483,7 @@ using core::Wavefront; hsa_status_t hsa_isa_from_name( const char *name, hsa_isa_t *isa) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(name); IS_BAD_PTR(isa); @@ -1553,7 +1495,7 @@ hsa_status_t hsa_isa_from_name( *isa = Isa::Handle(isa_object); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_agent_iterate_isas( @@ -1561,7 +1503,7 @@ hsa_status_t hsa_agent_iterate_isas( hsa_status_t (*callback)(hsa_isa_t isa, void *data), void *data) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(callback); @@ -1574,7 +1516,7 @@ hsa_status_t hsa_agent_iterate_isas( } return callback(Isa::Handle(isa_object), data); - CATCH + CATCH; } /* deprecated */ @@ -1583,7 +1525,7 @@ hsa_status_t hsa_isa_get_info( hsa_isa_info_t attribute, uint32_t index, void *value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -1596,14 +1538,14 @@ hsa_status_t hsa_isa_get_info( return isa_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; - CATCH + CATCH; } hsa_status_t hsa_isa_get_info_alt( hsa_isa_t isa, hsa_isa_info_t attribute, void *value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -1612,14 +1554,14 @@ hsa_status_t hsa_isa_get_info_alt( return isa_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; - CATCH + CATCH; } hsa_status_t hsa_isa_get_exception_policies( hsa_isa_t isa, hsa_profile_t profile, uint16_t *mask) { - TRY + TRY; IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_PTR(mask); @@ -1630,7 +1572,7 @@ hsa_status_t hsa_isa_get_exception_policies( // FIXME: update when exception policies are supported. *mask = 0; return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_isa_get_round_method( @@ -1638,7 +1580,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 + TRY; IS_OPEN(); IS_BAD_FP_TYPE(fp_type); IS_BAD_FLUSH_MODE(flush_mode); @@ -1649,14 +1591,14 @@ hsa_status_t hsa_isa_get_round_method( *round_method = isa_object->GetRoundMethod(fp_type, flush_mode); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_wavefront_get_info( hsa_wavefront_t wavefront, hsa_wavefront_info_t attribute, void *value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -1667,7 +1609,7 @@ hsa_status_t hsa_wavefront_get_info( return wavefront_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; - CATCH + CATCH; } hsa_status_t hsa_isa_iterate_wavefronts( @@ -1675,7 +1617,7 @@ hsa_status_t hsa_isa_iterate_wavefronts( hsa_status_t (*callback)(hsa_wavefront_t wavefront, void *data), void *data) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(callback); @@ -1686,7 +1628,7 @@ hsa_status_t hsa_isa_iterate_wavefronts( assert(wavefront_object); return callback(Wavefront::Handle(wavefront_object), data); - CATCH + CATCH; } /* deprecated */ @@ -1694,7 +1636,7 @@ hsa_status_t hsa_isa_compatible( hsa_isa_t code_object_isa, hsa_isa_t agent_isa, bool *result) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(result); @@ -1706,7 +1648,7 @@ hsa_status_t hsa_isa_compatible( *result = code_object_isa_object->IsCompatible(agent_isa_object); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } //===--- Code Objects (deprecated) ----------------------------------------===// @@ -1781,7 +1723,7 @@ hsa_status_t hsa_code_object_serialize( const char *options, void **serialized_code_object, size_t *serialized_code_object_size) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(alloc_callback); IS_BAD_PTR(serialized_code_object); @@ -1803,7 +1745,7 @@ hsa_status_t hsa_code_object_serialize( *serialized_code_object_size = code->ElfSize(); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } /* deprecated */ @@ -1812,7 +1754,7 @@ hsa_status_t hsa_code_object_deserialize( size_t serialized_code_object_size, const char *options, hsa_code_object_t *code_object) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(serialized_code_object); IS_BAD_PTR(code_object); @@ -1843,13 +1785,13 @@ hsa_status_t hsa_code_object_deserialize( code_object->handle = reinterpret_cast(code_object_alloc_data); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } /* deprecated */ hsa_status_t hsa_code_object_destroy( hsa_code_object_t code_object) { - TRY + TRY; IS_OPEN(); void *code_object_data = reinterpret_cast(code_object.handle); @@ -1863,7 +1805,7 @@ hsa_status_t hsa_code_object_destroy( HSA::hsa_memory_free(code_object_data); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } /* deprecated */ @@ -1871,7 +1813,7 @@ hsa_status_t hsa_code_object_get_info( hsa_code_object_t code_object, hsa_code_object_info_t attribute, void *value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -1901,7 +1843,7 @@ hsa_status_t hsa_code_object_get_info( return code->GetInfo(attribute, value); } } - CATCH + CATCH; } /* deprecated */ @@ -1909,7 +1851,7 @@ hsa_status_t hsa_code_object_get_symbol( hsa_code_object_t code_object, const char *symbol_name, hsa_code_symbol_t *symbol) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -1920,7 +1862,7 @@ hsa_status_t hsa_code_object_get_symbol( } return code->GetSymbol(nullptr, symbol_name, symbol); - CATCH + CATCH; } /* deprecated */ @@ -1929,7 +1871,7 @@ hsa_status_t hsa_code_object_get_symbol_from_name( const char *module_name, const char *symbol_name, hsa_code_symbol_t *symbol) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -1940,7 +1882,7 @@ hsa_status_t hsa_code_object_get_symbol_from_name( } return code->GetSymbol(module_name, symbol_name, symbol); - CATCH + CATCH; } /* deprecated */ @@ -1948,7 +1890,7 @@ hsa_status_t hsa_code_symbol_get_info( hsa_code_symbol_t code_symbol, hsa_code_symbol_info_t attribute, void *value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -1958,7 +1900,7 @@ hsa_status_t hsa_code_symbol_get_info( } return symbol->GetInfo(attribute, value); - CATCH + CATCH; } /* deprecated */ @@ -1968,7 +1910,7 @@ hsa_status_t hsa_code_object_iterate_symbols( hsa_code_symbol_t symbol, void *data), void *data) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(callback); @@ -1978,7 +1920,7 @@ hsa_status_t hsa_code_object_iterate_symbols( } return code->IterateSymbols(code_object, callback, data); - CATCH + CATCH; } //===--- Executable -------------------------------------------------------===// @@ -2032,7 +1974,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 + TRY; IS_OPEN(); IS_BAD_PTR(code_object_reader); @@ -2062,14 +2004,14 @@ hsa_status_t hsa_code_object_reader_create_from_file( *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper); return HSA_STATUS_SUCCESS; - CATCH + 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 + TRY; IS_OPEN(); IS_BAD_PTR(code_object); IS_BAD_PTR(code_object_reader); @@ -2084,12 +2026,12 @@ hsa_status_t hsa_code_object_reader_create_from_memory( *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_code_object_reader_destroy( hsa_code_object_reader_t code_object_reader) { - TRY + TRY; IS_OPEN(); CodeObjectReaderWrapper *wrapper = CodeObjectReaderWrapper::Object( @@ -2104,7 +2046,7 @@ hsa_status_t hsa_code_object_reader_destroy( delete wrapper; return HSA_STATUS_SUCCESS; - CATCH + CATCH; } /* deprecated */ @@ -2113,7 +2055,7 @@ hsa_status_t hsa_executable_create( hsa_executable_state_t executable_state, const char *options, hsa_executable_t *executable) { - TRY + TRY; IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_EXECUTABLE_STATE(executable_state); @@ -2136,7 +2078,7 @@ hsa_status_t hsa_executable_create( } return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_executable_create_alt( @@ -2144,7 +2086,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 + TRY; IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_ROUNDING_MODE(default_float_rounding_mode); // NOTES: should we check @@ -2159,12 +2101,12 @@ hsa_status_t hsa_executable_create_alt( *executable = Executable::Handle(exec); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_executable_destroy( hsa_executable_t executable) { - TRY + TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -2174,7 +2116,7 @@ hsa_status_t hsa_executable_destroy( GetLoader()->DestroyExecutable(exec); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } /* deprecated */ @@ -2183,7 +2125,7 @@ hsa_status_t hsa_executable_load_code_object( hsa_agent_t agent, hsa_code_object_t code_object, const char *options) { - TRY + TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -2192,7 +2134,7 @@ hsa_status_t hsa_executable_load_code_object( } return exec->LoadCodeObject(agent, code_object, options); - CATCH + CATCH; } hsa_status_t hsa_executable_load_program_code_object( @@ -2200,7 +2142,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 + TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -2218,7 +2160,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 + CATCH; } hsa_status_t hsa_executable_load_agent_code_object( @@ -2227,7 +2169,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 + TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -2245,13 +2187,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 + CATCH; } hsa_status_t hsa_executable_freeze( hsa_executable_t executable, const char *options) { - TRY + TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); @@ -2260,14 +2202,14 @@ hsa_status_t hsa_executable_freeze( } return exec->Freeze(options); - CATCH + CATCH; } hsa_status_t hsa_executable_get_info( hsa_executable_t executable, hsa_executable_info_t attribute, void *value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -2277,14 +2219,14 @@ hsa_status_t hsa_executable_get_info( } return exec->GetInfo(attribute, value); - CATCH + CATCH; } hsa_status_t hsa_executable_global_variable_define( hsa_executable_t executable, const char *variable_name, void *address) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(variable_name); @@ -2294,7 +2236,7 @@ hsa_status_t hsa_executable_global_variable_define( } return exec->DefineProgramExternalVariable(variable_name, address); - CATCH + CATCH; } hsa_status_t hsa_executable_agent_global_variable_define( @@ -2302,7 +2244,7 @@ hsa_status_t hsa_executable_agent_global_variable_define( hsa_agent_t agent, const char *variable_name, void *address) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(variable_name); @@ -2313,7 +2255,7 @@ hsa_status_t hsa_executable_agent_global_variable_define( return exec->DefineAgentExternalVariable( variable_name, agent, HSA_VARIABLE_SEGMENT_GLOBAL, address); - CATCH + CATCH; } hsa_status_t hsa_executable_readonly_variable_define( @@ -2321,7 +2263,7 @@ hsa_status_t hsa_executable_readonly_variable_define( hsa_agent_t agent, const char *variable_name, void *address) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(variable_name); @@ -2332,13 +2274,13 @@ hsa_status_t hsa_executable_readonly_variable_define( return exec->DefineAgentExternalVariable( variable_name, agent, HSA_VARIABLE_SEGMENT_READONLY, address); - CATCH + CATCH; } hsa_status_t hsa_executable_validate( hsa_executable_t executable, uint32_t *result) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(result); @@ -2348,19 +2290,19 @@ hsa_status_t hsa_executable_validate( } return exec->Validate(result); - CATCH + CATCH; } hsa_status_t hsa_executable_validate_alt( hsa_executable_t executable, const char *options, uint32_t *result) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(result); return HSA::hsa_executable_validate(executable, result); - CATCH + CATCH; } /* deprecated */ @@ -2371,7 +2313,7 @@ hsa_status_t hsa_executable_get_symbol( hsa_agent_t agent, int32_t call_convention, hsa_executable_symbol_t *symbol) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -2394,7 +2336,7 @@ hsa_status_t hsa_executable_get_symbol( return HSA::hsa_executable_get_symbol_by_name( executable, mangled_name.c_str(), exec->IsProgramSymbol(mangled_name.c_str()) ? nullptr : &agent, symbol); - CATCH + CATCH; } hsa_status_t hsa_executable_get_symbol_by_name( @@ -2404,7 +2346,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 + TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); @@ -2421,14 +2363,14 @@ hsa_status_t hsa_executable_get_symbol_by_name( *symbol = loader::Symbol::Handle(sym); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_executable_symbol_get_info( hsa_executable_symbol_t executable_symbol, hsa_executable_symbol_info_t attribute, void *value) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -2439,7 +2381,7 @@ hsa_status_t hsa_executable_symbol_get_info( return sym->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; - CATCH + CATCH; } /* deprecated */ @@ -2449,7 +2391,7 @@ hsa_status_t hsa_executable_iterate_symbols( hsa_executable_symbol_t symbol, void *data), void *data) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(callback); @@ -2459,7 +2401,7 @@ hsa_status_t hsa_executable_iterate_symbols( } return exec->IterateSymbols(callback, data); - CATCH + CATCH; } hsa_status_t hsa_executable_iterate_agent_symbols( @@ -2470,7 +2412,7 @@ hsa_status_t hsa_executable_iterate_agent_symbols( hsa_executable_symbol_t symbol, void *data), void *data) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(callback); @@ -2484,7 +2426,7 @@ hsa_status_t hsa_executable_iterate_agent_symbols( } return exec->IterateAgentSymbols(agent, callback, data); - CATCH + CATCH; } hsa_status_t hsa_executable_iterate_program_symbols( @@ -2493,7 +2435,7 @@ hsa_status_t hsa_executable_iterate_program_symbols( hsa_executable_symbol_t symbol, void *data), void *data) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(callback); @@ -2503,7 +2445,7 @@ hsa_status_t hsa_executable_iterate_program_symbols( } return exec->IterateProgramSymbols(callback, data); - CATCH + CATCH; } //===--- Runtime Notifications --------------------------------------------===// diff --git a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp index 7979ea13ab..75814a8d16 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp @@ -378,6 +378,9 @@ void HsaApiTable::UpdateAmdExts() { amd_ext_api.hsa_amd_ipc_memory_create_fn = AMD::hsa_amd_ipc_memory_create; amd_ext_api.hsa_amd_ipc_memory_attach_fn = AMD::hsa_amd_ipc_memory_attach; amd_ext_api.hsa_amd_ipc_memory_detach_fn = AMD::hsa_amd_ipc_memory_detach; + amd_ext_api.hsa_amd_signal_create_fn = AMD::hsa_amd_signal_create; + amd_ext_api.hsa_amd_ipc_signal_create_fn = AMD::hsa_amd_ipc_signal_create; + amd_ext_api.hsa_amd_ipc_signal_attach_fn = AMD::hsa_amd_ipc_signal_attach; } class Init { diff --git a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp index 27f7f537bc..a0f63c7b1a 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp @@ -43,6 +43,7 @@ #include #include #include +#include #include "hsakmt.h" @@ -52,7 +53,10 @@ #include "core/inc/amd_gpu_agent.h" #include "core/inc/amd_memory_region.h" #include "core/inc/signal.h" +#include "core/inc/default_signal.h" #include "core/inc/interrupt_signal.h" +#include "core/inc/ipc_signal.h" +#include "core/inc/exceptions.h" template struct ValidityError; @@ -124,11 +128,13 @@ hsa_status_t handleException() { 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 (const hsa_exception& e) { + return e.error_code(); + // Enable when callback exception support is added. + // } 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(); @@ -143,7 +149,7 @@ template static __forceinline T handleExceptionT() { } hsa_status_t hsa_amd_coherency_get_type(hsa_agent_t agent_handle, hsa_amd_coherency_type_t* type) { - TRY + TRY; IS_OPEN(); const core::Agent* agent = core::Agent::Convert(agent_handle); @@ -162,12 +168,12 @@ hsa_status_t hsa_amd_coherency_get_type(hsa_agent_t agent_handle, hsa_amd_cohere *type = gpu_agent->current_coherency_type(); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle, hsa_amd_coherency_type_t type) { - TRY + TRY; IS_OPEN(); core::Agent* agent = core::Agent::Convert(agent_handle); @@ -190,11 +196,11 @@ hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle, } return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) { - TRY + TRY; IS_OPEN(); if (ptr == NULL) { @@ -206,14 +212,14 @@ hsa_status_t hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count) { } return core::Runtime::runtime_singleton_->FillMemory(ptr, value, count); - CATCH + 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) { - TRY + TRY; if (dst == NULL || src == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } if ((num_dep_signals == 0 && dep_signals != NULL) || @@ -246,11 +252,11 @@ hsa_status_t hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent_handle, } return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable) { - TRY + TRY; IS_OPEN(); core::Queue* cmd_queue = core::Queue::Convert(queue); @@ -261,11 +267,11 @@ hsa_status_t hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enab AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, (enable != 0)); return HSA_STATUS_SUCCESS; - CATCH + CATCH; } hsa_status_t hsa_amd_profiling_async_copy_enable(bool enable) { - TRY + TRY; IS_OPEN(); return core::Runtime::runtime_singleton_->IterateAgent( @@ -274,13 +280,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 + 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 + TRY; IS_OPEN(); IS_BAD_PTR(time); @@ -303,12 +309,12 @@ hsa_status_t hsa_amd_profiling_get_dispatch_time( gpu_agent->TranslateTime(signal, *time); return HSA_STATUS_SUCCESS; - CATCH + 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 + TRY; IS_OPEN(); IS_BAD_PTR(time); @@ -333,13 +339,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 + 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) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(system_tick); @@ -357,42 +363,89 @@ hsa_status_t hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent_h *system_tick = gpu_agent->TranslateTime(agent_tick); return HSA_STATUS_SUCCESS; - CATCH + CATCH; +} + +hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, + const hsa_agent_t* consumers, uint64_t attributes, + hsa_signal_t* hsa_signal) { + struct AgentHandleCompare { + bool operator()(const hsa_agent_t& lhs, const hsa_agent_t& rhs) const { + return lhs.handle < rhs.handle; + } + }; + + TRY; + IS_OPEN(); + IS_BAD_PTR(hsa_signal); + + core::Signal* ret; + + bool enable_ipc = attributes & HSA_AMD_SIGNAL_IPC; + bool use_default = + enable_ipc || (attributes & HSA_AMD_SIGNAL_AMD_GPU_ONLY) || (!core::g_use_interrupt_wait); + + if ((!use_default) && (num_consumers != 0)) { + IS_BAD_PTR(consumers); + + // Check for duplicates in consumers. + std::set consumer_set(consumers, consumers + num_consumers); + if (consumer_set.size() != num_consumers) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + use_default = true; + for (const core::Agent* cpu_agent : core::Runtime::runtime_singleton_->cpu_agents()) { + use_default &= (consumer_set.find(cpu_agent->public_handle()) == consumer_set.end()); + } + } + + if (use_default) { + ret = new core::DefaultSignal(initial_value, enable_ipc); + } else { + ret = new core::InterruptSignal(initial_value); + } + + *hsa_signal = core::Signal::Convert(ret); + 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) { - TRY + TRY; + IS_OPEN(); // 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(core::SharedSignal::Convert(hsa_signals[i])->IsValid() && "Invalid signal."); + assert(hsa_signals[i].handle != 0 && 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) + 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) { - TRY + TRY; IS_OPEN(); + IS_BAD_PTR(handler); core::Signal* signal = core::Signal::Convert(hsa_signal); IS_VALID(signal); - IS_BAD_PTR(handler); if (core::g_use_interrupt_wait && (!core::InterruptSignal::IsType(signal))) return HSA_STATUS_ERROR_INVALID_SIGNAL; return core::Runtime::runtime_singleton_->SetAsyncSignalHandler( hsa_signal, cond, value, handler, arg); - CATCH + CATCH; } hsa_status_t hsa_amd_async_function(void (*callback)(void* arg), void* arg) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(callback); @@ -400,26 +453,26 @@ hsa_status_t hsa_amd_async_function(void (*callback)(void* arg), void* arg) { return core::Runtime::runtime_singleton_->SetAsyncSignalHandler( null_signal, HSA_SIGNAL_CONDITION_EQ, 0, (hsa_amd_signal_handler)callback, arg); - CATCH + 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 + 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 + CATCH; } hsa_status_t hsa_amd_memory_lock(void* host_ptr, size_t size, hsa_agent_t* agents, int num_agent, void** agent_ptr) { - TRY + TRY; IS_OPEN(); *agent_ptr = NULL; @@ -437,11 +490,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 + CATCH; } hsa_status_t hsa_amd_memory_unlock(void* host_ptr) { - TRY + TRY; IS_OPEN(); const amd::MemoryRegion* system_region = @@ -449,12 +502,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 + 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) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(value); @@ -465,14 +518,14 @@ hsa_status_t hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool, } return mem_region->GetPoolInfo(attribute, value); - CATCH + 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 + TRY; IS_OPEN(); IS_BAD_PTR(callback); const core::Agent* agent = core::Agent::Convert(agent_handle); @@ -490,12 +543,12 @@ hsa_status_t hsa_amd_agent_iterate_memory_pools( reinterpret_cast( callback), data); - CATCH + CATCH; } hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size, uint32_t flags, void** ptr) { - TRY + TRY; IS_OPEN(); if (size == 0 || ptr == NULL) { @@ -511,7 +564,7 @@ hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, siz return core::Runtime::runtime_singleton_->AllocateMemory( mem_region, size, core::MemoryRegion::AllocateRestrict, ptr); - CATCH + CATCH; } hsa_status_t hsa_amd_memory_pool_free(void* ptr) { @@ -520,7 +573,7 @@ hsa_status_t hsa_amd_memory_pool_free(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 + TRY; IS_OPEN(); if (num_agents == 0 || agents == NULL || flags != NULL || ptr == NULL) { @@ -529,12 +582,12 @@ hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* return core::Runtime::runtime_singleton_->AllowAccess(num_agents, agents, ptr); - CATCH + 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) { - TRY + TRY; IS_OPEN(); if (result == NULL) { @@ -558,13 +611,13 @@ hsa_status_t hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_po } return src_mem_region->CanMigrate(*dst_mem_region, *result); - CATCH + CATCH; } hsa_status_t hsa_amd_memory_migrate(const void* ptr, hsa_amd_memory_pool_t memory_pool, uint32_t flags) { - TRY + TRY; IS_OPEN(); if (ptr == NULL || flags != 0) { @@ -580,13 +633,13 @@ hsa_status_t hsa_amd_memory_migrate(const void* ptr, } return dst_mem_region->Migrate(flags, ptr); - CATCH + 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 + TRY; IS_OPEN(); if (value == NULL) { @@ -605,7 +658,7 @@ hsa_status_t hsa_amd_agent_memory_pool_get_info( } return mem_region->GetAgentPoolInfo(*agent, attribute, value); - CATCH + CATCH; } hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents, @@ -614,7 +667,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 + TRY; IS_OPEN(); IS_BAD_PTR(agents); IS_BAD_PTR(size); @@ -641,49 +694,49 @@ hsa_status_t hsa_amd_interop_map_buffer(uint32_t num_agents, if (num_agents > tinyArraySize) delete[] core_agents; return ret; - CATCH + CATCH; } hsa_status_t hsa_amd_interop_unmap_buffer(void* ptr) { - TRY + TRY; IS_OPEN(); if (ptr != NULL) core::Runtime::runtime_singleton_->InteropUnmap(ptr); return HSA_STATUS_SUCCESS; - CATCH + 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) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(ptr); IS_BAD_PTR(info); return core::Runtime::runtime_singleton_->PtrInfo(ptr, info, alloc, num_accessible, accessible); - CATCH + CATCH; } hsa_status_t hsa_amd_pointer_info_set_userdata(void* ptr, void* userdata) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(ptr); return core::Runtime::runtime_singleton_->SetPtrInfoData(ptr, userdata); - CATCH + CATCH; } hsa_status_t hsa_amd_ipc_memory_create(void* ptr, size_t len, hsa_amd_ipc_memory_t* handle) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(ptr); IS_BAD_PTR(handle); return core::Runtime::runtime_singleton_->IPCCreate(ptr, len, handle); - CATCH + 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 + TRY; IS_OPEN(); IS_BAD_PTR(mapped_ptr); if (num_agents != 0) IS_BAD_PTR(mapping_agents); @@ -706,15 +759,38 @@ 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 + CATCH; } hsa_status_t hsa_amd_ipc_memory_detach(void* mapped_ptr) { - TRY + TRY; IS_OPEN(); IS_BAD_PTR(mapped_ptr); return core::Runtime::runtime_singleton_->IPCDetach(mapped_ptr); - CATCH + CATCH; +} + +hsa_status_t hsa_amd_ipc_signal_create(hsa_signal_t hsa_signal, hsa_amd_ipc_signal_t* handle) { + TRY; + IS_OPEN(); + IS_BAD_PTR(handle); + core::Signal* signal = core::Signal::Convert(hsa_signal); + IS_VALID(signal); + core::IPCSignal::CreateHandle(signal, handle); + return HSA_STATUS_SUCCESS; + CATCH; +} + +hsa_status_t hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle, + hsa_signal_t* hsa_signal) { + TRY; + IS_OPEN(); + IS_BAD_PTR(handle); + IS_BAD_PTR(hsa_signal); + core::Signal* signal = core::IPCSignal::Attach(handle); + *hsa_signal = core::Signal::Convert(signal); + return HSA_STATUS_SUCCESS; + CATCH; } } // end of AMD namespace diff --git a/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp b/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp index 7f80164996..c22a906044 100644 --- a/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp +++ b/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp @@ -89,10 +89,6 @@ InterruptSignal::InterruptSignal(hsa_signal_value_t initial_value, HsaEvent* use } InterruptSignal::~InterruptSignal() { - invalid_ = true; - SetEvent(); - while (InUse()) - ; if (free_event_) hsaKmtDestroyEvent(event_); } @@ -119,14 +115,14 @@ void InterruptSignal::StoreRelease(hsa_signal_value_t value) { hsa_signal_value_t InterruptSignal::WaitRelaxed( hsa_signal_condition_t condition, hsa_signal_value_t compare_value, uint64_t timeout, hsa_wait_state_t wait_hint) { - uint32_t prior = atomic::Increment(&waiting_); + Retain(); + MAKE_SCOPE_GUARD([&]() { Release(); }); - // assert(prior == 0 && "Multiple waiters on interrupt signal!"); + uint32_t prior = waiting_++; + MAKE_SCOPE_GUARD([&]() { waiting_--; }); // Allow only the first waiter to sleep (temporary, known to be bad). if (prior != 0) wait_hint = HSA_WAIT_STATE_ACTIVE; - MAKE_SCOPE_GUARD([&]() { atomic::Decrement(&waiting_); }); - int64_t value; timer::fast_clock::time_point start_time = timer::fast_clock::now(); @@ -143,7 +139,7 @@ hsa_signal_value_t InterruptSignal::WaitRelaxed( bool condition_met = false; while (true) { - if (invalid_) return 0; + if (!IsValid()) return 0; value = atomic::Load(&signal_.value, std::memory_order_relaxed); diff --git a/runtime/hsa-runtime/core/runtime/ipc_signal.cpp b/runtime/hsa-runtime/core/runtime/ipc_signal.cpp new file mode 100644 index 0000000000..4b4d62fef3 --- /dev/null +++ b/runtime/hsa-runtime/core/runtime/ipc_signal.cpp @@ -0,0 +1,94 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// The University of Illinois/NCSA +// Open Source License (NCSA) +// +// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +// +// Developed by: +// +// AMD Research and AMD HSA Software Development +// +// Advanced Micro Devices, Inc. +// +// www.amd.com +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to +// deal with the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimers. +// - Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimers in +// the documentation and/or other materials provided with the distribution. +// - Neither the names of Advanced Micro Devices, Inc, +// nor the names of its contributors may be used to endorse or promote +// products derived from this Software without specific prior written +// permission. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +// DEALINGS WITH THE SOFTWARE. +// +//////////////////////////////////////////////////////////////////////////////// + +#include "core/inc/ipc_signal.h" + +#include + +#include "core/inc/runtime.h" +#include "core/inc/exceptions.h" + +namespace core { + +int IPCSignal::rtti_id_ = 0; +KernelMutex IPCSignal::lock_; + +SharedMemory::SharedMemory(const hsa_amd_ipc_memory_t* handle, size_t len) { + hsa_status_t err = Runtime::runtime_singleton_->IPCAttach(handle, len, 0, NULL, &ptr_); + if (err != HSA_STATUS_SUCCESS) throw AMD::hsa_exception(err, "IPC memory attach failed."); +} + +SharedMemory::SharedMemory(SharedMemory&& rhs) { + ptr_ = rhs.ptr_; + rhs.ptr_ = nullptr; +} + +SharedMemory::~SharedMemory() { + if (ptr_ == nullptr) return; + auto err = Runtime::runtime_singleton_->IPCDetach(ptr_); + assert(err == HSA_STATUS_SUCCESS && "IPC detach failed."); +} + +void IPCSignal::CreateHandle(Signal* signal, hsa_amd_ipc_signal_t* ipc_handle) { + if (!signal->isIPC()) + throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT, "Signal must be IPC enabled."); + SharedSignal* shared = SharedSignal::Convert(Convert(signal)); + hsa_status_t err = Runtime::runtime_singleton_->IPCCreate(shared, 4096, ipc_handle); + if (err != HSA_STATUS_SUCCESS) throw AMD::hsa_exception(err, "IPC memory create failed."); +} + +Signal* IPCSignal::Attach(const hsa_amd_ipc_signal_t* ipc_signal_handle) { + SharedMemorySignal shared(ipc_signal_handle); + + if (!(shared.signal()->IsIPC())) + throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_ARGUMENT, + "IPC memory does not contain an IPC signal abi block."); + + hsa_signal_t handle = SharedSignal::Convert(shared.signal()); + + ScopedAcquire lock(&lock_); + Signal* ret = core::Signal::DuplicateHandle(handle); + if (ret == nullptr) ret = new IPCSignal(std::move(shared)); + return ret; +} + +} // namespace core diff --git a/runtime/hsa-runtime/core/runtime/runtime.cpp b/runtime/hsa-runtime/core/runtime/runtime.cpp index b0e9d41a55..e6a28e86d3 100644 --- a/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -1026,7 +1026,7 @@ void Runtime::Unload() { async_events_control_.Shutdown(); - delete vm_fault_signal_; + vm_fault_signal_->DestroySignal(); core::InterruptSignal::DestroyEvent(vm_fault_event_); DestroyAgents(); diff --git a/runtime/hsa-runtime/core/runtime/signal.cpp b/runtime/hsa-runtime/core/runtime/signal.cpp index 8b26396849..5400e1be18 100644 --- a/runtime/hsa-runtime/core/runtime/signal.cpp +++ b/runtime/hsa-runtime/core/runtime/signal.cpp @@ -44,24 +44,83 @@ #define HSA_RUNTME_CORE_SIGNAL_CPP_ #include "core/inc/signal.h" -#include "core/util/timer.h" + #include +#include "core/util/timer.h" namespace core { +KernelMutex Signal::ipcLock_; +std::map Signal::ipcMap_; + +void Signal::registerIpc() { + ScopedAcquire lock(&ipcLock_); + auto handle = Convert(this); + assert(ipcMap_.find(handle.handle) == ipcMap_.end() && + "Can't register the same IPC signal twice."); + ipcMap_[handle.handle] = this; +} + +bool Signal::deregisterIpc() { + ScopedAcquire lock(&ipcLock_); + if (refcount_ != 0) return false; + auto handle = Convert(this); + const auto& it = ipcMap_.find(handle.handle); + assert(it != ipcMap_.end() && "Deregister on non-IPC signal."); + ipcMap_.erase(it); + return true; +} + +Signal* Signal::lookupIpc(hsa_signal_t signal) { + ScopedAcquire lock(&ipcLock_); + const auto& it = ipcMap_.find(signal.handle); + if (it == ipcMap_.end()) return nullptr; + return it->second; +} + +Signal* Signal::duplicateIpc(hsa_signal_t signal) { + ScopedAcquire lock(&ipcLock_); + const auto& it = ipcMap_.find(signal.handle); + if (it == ipcMap_.end()) return nullptr; + it->second->refcount_++; + it->second->Retain(); + return it->second; +} + +void Signal::Release() { + if (--retained_ != 0) return; + if (!isIPC()) + doDestroySignal(); + else if (deregisterIpc()) + doDestroySignal(); +} + +Signal::~Signal() { + signal_.kind = AMD_SIGNAL_KIND_INVALID; + if (refcount_ == 1 && isIPC()) { + refcount_ = 0; + deregisterIpc(); + } +} + uint32_t Signal::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, hsa_wait_state_t wait_hint, hsa_signal_value_t* satisfying_value) { hsa_signal_handle* signals = reinterpret_cast(const_cast(hsa_signals)); - uint32_t prior = 0; - for (uint32_t i = 0; i < signal_count; i++) - prior = Max(prior, atomic::Increment(&signals[i]->waiting_)); + + for (uint32_t i = 0; i < signal_count; i++) signals[i]->Retain(); MAKE_SCOPE_GUARD([&]() { - for (uint32_t i = 0; i < signal_count; i++) - atomic::Decrement(&signals[i]->waiting_); + for (uint32_t i = 0; i < signal_count; i++) signals[i]->Release(); + }); + + uint32_t prior = 0; + for (uint32_t i = 0; i < signal_count; i++) prior = Max(prior, signals[i]->waiting_++); + + MAKE_SCOPE_GUARD([&]() { + for (uint32_t i = 0; i < signal_count; i++) signals[i]->waiting_--; }); // Allow only the first waiter to sleep (temporary, known to be bad). @@ -113,7 +172,7 @@ uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals, bool condition_met = false; while (true) { for (uint32_t i = 0; i < signal_count; i++) { - if (signals[i]->invalid_) return uint32_t(-1); + if (!signals[i]->IsValid()) return uint32_t(-1); // Handling special event. if (signals[i]->EopEvent() != NULL) { diff --git a/runtime/hsa-runtime/hsacore.so.def b/runtime/hsa-runtime/hsacore.so.def index d3bc264289..e953304d64 100644 --- a/runtime/hsa-runtime/hsacore.so.def +++ b/runtime/hsa-runtime/hsacore.so.def @@ -173,6 +173,7 @@ global: hsa_amd_profiling_async_copy_enable; hsa_amd_profiling_get_async_copy_time; hsa_amd_profiling_convert_tick_to_system_domain; + hsa_amd_signal_create; hsa_amd_signal_wait_any; hsa_amd_signal_async_handler; hsa_amd_async_function; @@ -203,14 +204,16 @@ global: hsa_ext_image_destroy; hsa_ext_sampler_create; hsa_ext_sampler_destroy; - hsa_ext_image_get_capability_with_layout; - hsa_ext_image_data_get_info_with_layout; - hsa_ext_image_create_with_layout; + hsa_ext_image_get_capability_with_layout; + hsa_ext_image_data_get_info_with_layout; + hsa_ext_image_create_with_layout; hsa_amd_pointer_info; hsa_amd_pointer_info_set_userdata; hsa_amd_ipc_memory_create; hsa_amd_ipc_memory_attach; hsa_amd_ipc_memory_detach; + hsa_amd_ipc_signal_create; + hsa_amd_ipc_signal_attach; local: *; diff --git a/runtime/hsa-runtime/inc/hsa_api_trace.h b/runtime/hsa-runtime/inc/hsa_api_trace.h index 3bbb0608a9..a1927198d2 100644 --- a/runtime/hsa-runtime/inc/hsa_api_trace.h +++ b/runtime/hsa-runtime/inc/hsa_api_trace.h @@ -166,6 +166,9 @@ struct AmdExtTable { decltype(hsa_amd_ipc_memory_create)* hsa_amd_ipc_memory_create_fn; decltype(hsa_amd_ipc_memory_attach)* hsa_amd_ipc_memory_attach_fn; decltype(hsa_amd_ipc_memory_detach)* hsa_amd_ipc_memory_detach_fn; + decltype(hsa_amd_signal_create)* hsa_amd_signal_create_fn; + decltype(hsa_amd_ipc_signal_create)* hsa_amd_ipc_signal_create_fn; + decltype(hsa_amd_ipc_signal_attach)* hsa_amd_ipc_signal_attach_fn; }; // Table to export HSA Core Runtime Apis diff --git a/runtime/hsa-runtime/inc/hsa_ext_amd.h b/runtime/hsa-runtime/inc/hsa_ext_amd.h index bc026f26cc..d4a1841240 100755 --- a/runtime/hsa-runtime/inc/hsa_ext_amd.h +++ b/runtime/hsa-runtime/inc/hsa_ext_amd.h @@ -377,6 +377,62 @@ hsa_status_t HSA_API uint64_t agent_tick, uint64_t* system_tick); +/** + * @brief Signal attribute flags. + */ +typedef enum { + /** + * Signal will only be consumed by AMD GPUs. Limits signal consumption to + * AMD GPU agents only. Ignored if @p num_consumers is not zero (all agents). + */ + HSA_AMD_SIGNAL_AMD_GPU_ONLY = 1, + /** + * Signal may be used for interprocess communication. + * This signal may not be used with profiling APIs. Errors or inaccurate + * timing data may result from such use. + */ + HSA_AMD_SIGNAL_IPC = 2, +} hsa_amd_signal_attribute_t; + +/** + * @brief Create a signal with specific attributes. + * + * @param[in] initial_value Initial value of the signal. + * + * @param[in] num_consumers Size of @p consumers. A value of 0 indicates that + * any agent might wait on the signal. + * + * @param[in] consumers List of agents that might consume (wait on) the + * signal. If @p num_consumers is 0, this argument is ignored; otherwise, the + * HSA runtime might use the list to optimize the handling of the signal + * object. If an agent not listed in @p consumers waits on the returned + * signal, the behavior is undefined. The memory associated with @p consumers + * can be reused or freed after the function returns. + * + * @param[in] attributes Requested signal attributes. Multiple signal attributes + * may be requested by combining them with bitwise OR. Requesting no attributes + * (@p attributes == 0) results in the same signal as would have been obtained + * via hsa_signal_create. + * + * @param[out] signal Pointer to a memory location where the HSA runtime will + * store the newly created signal handle. Must not be NULL. + * + * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully. + * + * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been + * initialized. + * + * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate + * the required resources. + * + * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p signal is NULL, @p + * num_consumers is greater than 0 but @p consumers is NULL, or @p consumers + * contains duplicates. + */ +hsa_status_t HSA_API hsa_amd_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, + const hsa_agent_t* consumers, uint64_t attributes, + hsa_signal_t* signal); + /** * @brief Asyncronous signal handler function type. * @@ -1511,6 +1567,60 @@ hsa_status_t HSA_API hsa_amd_ipc_memory_attach( */ hsa_status_t HSA_API hsa_amd_ipc_memory_detach(void* mapped_ptr); +/** + * @brief 256-bit process independent identifier for a ROCr IPC signal. + */ +typedef hsa_amd_ipc_memory_t hsa_amd_ipc_signal_t; + +/** + * @brief Obtains an interprocess sharing handle for a signal. The handle is + * valid while the signal it references remains valid in any process. In + * general applications should confirm that the signal has been attached (via + * hsa_amd_ipc_signal_attach) in the remote process prior to destroying that + * signal in the local process. + * Repeated calls for the same signal may, but are not required to, return + * unique handles. + * + * @param[in] signal Signal created with attribute HSA_AMD_SIGNAL_IPC. + * + * @param[out] handle Process independent identifier referencing the shared + * signal. + * + * @retval HSA_STATUS_SUCCESS @p handle is ready to use for interprocess sharing. + * + * @retval HSA_STATUS_ERROR_NOT_INITIALIZED if HSA is not initialized + * + * @retval HSA_STATUS_ERROR_OUT_OF_RESOURCES if there is a failure in allocating + * necessary resources + * + * @retval HSA_STATUS_ERROR_INVALID_ARGUMENT @p signal is not a valid signal + * created with attribute HSA_AMD_SIGNAL_IPC or handle is NULL. + */ +hsa_status_t HSA_API hsa_amd_ipc_signal_create(hsa_signal_t signal, hsa_amd_ipc_signal_t* handle); + +/** + * @brief Imports an IPC capable signal into the local process. If an IPC + * signal handle is attached multiple times in a process each attach may return + * a different signal handle. Each returned signal handle is refcounted and + * requires a matching number of calls to hsa_signal_destroy to release the + * shared signal. + * + * @param[in] handle Pointer to the identifier for the shared signal. + * + * @param[out] signal Recieves a process local signal handle to the shared signal. + * + * @retval HSA_STATUS_SUCCESS if the signal is successfully imported. + * + * @retval HSA_STATUS_ERROR_NOT_INITIALIZED if HSA is not initialized + * + * @retval HSA_STATUS_ERROR_OUT_OF_RESOURCES if there is a failure in allocating + * necessary resources + * + * @retval HSA_STATUS_ERROR_INVALID_ARGUMENT @p handle is not valid. + */ +hsa_status_t HSA_API hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle, + hsa_signal_t* signal); + #ifdef __cplusplus } // end extern "C" block #endif