From c9642cf7af52cc1c45486787190cdb8bc391d485 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Wed, 5 Jul 2017 22:27:21 -0500 Subject: [PATCH] Initial IPC signal support. Added an API for creating signals with attributes. Added two APIs for IPC operations on signals. Initial use of exceptions for error handling. Add ref counting to signals. Removed spin loops from signal destructors. Signals are no longer to be destroyed with delete, use DeleteSignal instead. Added delete safety to doorbells. Added secondary hsa_signal_t -> Signal* translation path for IPC enabled signals. Change-Id: Id59065d002f0c2566b0a9425694da2ed27cb7d7f --- runtime/hsa-runtime/CMakeLists.txt | 1 + .../core/common/hsa_table_interface.cpp | 19 + runtime/hsa-runtime/core/inc/amd_aql_queue.h | 3 + runtime/hsa-runtime/core/inc/default_signal.h | 34 +- runtime/hsa-runtime/core/inc/exceptions.h | 67 +++ .../hsa-runtime/core/inc/hsa_ext_amd_impl.h | 11 + runtime/hsa-runtime/core/inc/ipc_signal.h | 112 ++++ runtime/hsa-runtime/core/inc/signal.h | 127 ++-- .../core/runtime/amd_gpu_agent.cpp | 6 +- .../core/runtime/default_signal.cpp | 111 ++-- runtime/hsa-runtime/core/runtime/hsa.cpp | 546 ++++++++---------- .../core/runtime/hsa_api_trace.cpp | 3 + .../hsa-runtime/core/runtime/hsa_ext_amd.cpp | 206 ++++--- .../core/runtime/interrupt_signal.cpp | 14 +- .../hsa-runtime/core/runtime/ipc_signal.cpp | 94 +++ runtime/hsa-runtime/core/runtime/runtime.cpp | 2 +- runtime/hsa-runtime/core/runtime/signal.cpp | 73 ++- runtime/hsa-runtime/hsacore.so.def | 9 +- runtime/hsa-runtime/inc/hsa_api_trace.h | 3 + runtime/hsa-runtime/inc/hsa_ext_amd.h | 110 ++++ 20 files changed, 1063 insertions(+), 488 deletions(-) create mode 100644 runtime/hsa-runtime/core/inc/exceptions.h create mode 100644 runtime/hsa-runtime/core/inc/ipc_signal.h create mode 100644 runtime/hsa-runtime/core/runtime/ipc_signal.cpp 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