From 0c7dde2d1f8ab04666fe9cd86011d03ba0e8b215 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Tue, 17 Oct 2017 22:09:12 -0500 Subject: [PATCH] Add queue intercept support to the runtime. Queue intercept is exposed as two tools-only APIs via the API intercept table. Change-Id: Iac9602ed3143974d85c3569e9092295ad18037f8 --- runtime/hsa-runtime/CMakeLists.txt | 1 + .../core/common/hsa_table_interface.cpp | 16 + runtime/hsa-runtime/core/common/shared.h | 5 +- runtime/hsa-runtime/core/inc/amd_aql_queue.h | 4 +- runtime/hsa-runtime/core/inc/host_queue.h | 6 + .../hsa-runtime/core/inc/intercept_queue.h | 413 ++++++++++++++++++ runtime/hsa-runtime/core/inc/queue.h | 6 + runtime/hsa-runtime/core/inc/runtime.h | 1 + .../core/runtime/amd_gpu_agent.cpp | 8 +- .../hsa-runtime/core/runtime/host_queue.cpp | 1 + .../core/runtime/hsa_api_trace.cpp | 13 + .../hsa-runtime/core/runtime/hsa_ext_amd.cpp | 48 +- .../core/runtime/intercept_queue.cpp | 260 +++++++++++ runtime/hsa-runtime/core/runtime/runtime.cpp | 4 +- runtime/hsa-runtime/core/util/utils.h | 18 +- runtime/hsa-runtime/inc/hsa_api_trace.h | 15 + 16 files changed, 800 insertions(+), 19 deletions(-) create mode 100644 runtime/hsa-runtime/core/inc/intercept_queue.h create mode 100644 runtime/hsa-runtime/core/runtime/intercept_queue.cpp diff --git a/runtime/hsa-runtime/CMakeLists.txt b/runtime/hsa-runtime/CMakeLists.txt index 5ca48bab8f..4240299fe5 100644 --- a/runtime/hsa-runtime/CMakeLists.txt +++ b/runtime/hsa-runtime/CMakeLists.txt @@ -135,6 +135,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/intercept_queue.cpp" "core/runtime/ipc_signal.cpp" "core/runtime/isa.cpp" "core/runtime/runtime.cpp" diff --git a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp b/runtime/hsa-runtime/core/common/hsa_table_interface.cpp index 5cc20378b7..85da28adce 100644 --- a/runtime/hsa-runtime/core/common/hsa_table_interface.cpp +++ b/runtime/hsa-runtime/core/common/hsa_table_interface.cpp @@ -1102,3 +1102,19 @@ hsa_status_t HSA_API hsa_amd_register_system_event_handler( void* data) { return amdExtTable->hsa_amd_register_system_event_handler_fn(type, callback, data); } + +// Mirrors Amd Extension Apis +hsa_status_t hsa_amd_queue_intercept_create( + hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type, + 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) { + return amdExtTable->hsa_amd_queue_intercept_create_fn( + agent_handle, size, type, callback, data, private_segment_size, group_segment_size, queue); +} + +// Mirrors Amd Extension Apis +hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue, + hsa_amd_queue_intercept_handler callback, + void* user_data) { + return amdExtTable->hsa_amd_queue_intercept_register_fn(queue, callback, user_data); +} diff --git a/runtime/hsa-runtime/core/common/shared.h b/runtime/hsa-runtime/core/common/shared.h index b8aced5e56..5ea97932d6 100644 --- a/runtime/hsa-runtime/core/common/shared.h +++ b/runtime/hsa-runtime/core/common/shared.h @@ -43,11 +43,12 @@ #ifndef HSA_RUNTME_CORE_INC_SHARED_H_ #define HSA_RUNTME_CORE_INC_SHARED_H_ -#include "core/util/utils.h" #include - #include #include +#include + +#include "core/util/utils.h" namespace core { /// @brief Base class encapsulating the allocator and deallocator for diff --git a/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/runtime/hsa-runtime/core/inc/amd_aql_queue.h index 0213a02b4b..4033ba61cd 100644 --- a/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -59,6 +59,8 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig return signal->IsType(&rtti_id_); } + static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); } + // Acquires/releases queue resources and requests HW schedule/deschedule. AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, ScratchInfo& scratch, core::HsaEventCallback callback, @@ -337,7 +339,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig } protected: - bool _IsA(rtti_t id) const override { return id == &rtti_id_; } + bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id_; } /// @brief Disallow destroying doorbell apart from its queue. void doDestroySignal() override { assert(false); } diff --git a/runtime/hsa-runtime/core/inc/host_queue.h b/runtime/hsa-runtime/core/inc/host_queue.h index 2eb1e25f4c..5a883ed327 100644 --- a/runtime/hsa-runtime/core/inc/host_queue.h +++ b/runtime/hsa-runtime/core/inc/host_queue.h @@ -51,6 +51,8 @@ namespace core { class HostQueue : public Queue { public: + static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); } + HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t type, uint32_t features, hsa_signal_t doorbell_signal); @@ -158,7 +160,11 @@ class HostQueue : public Queue { void operator delete(void*, void*) {} + protected: + bool _IsA(Queue::rtti_t id) const { return id == &rtti_id_; } + private: + static int rtti_id_; static const size_t kRingAlignment = 256; const uint32_t size_; bool active_; diff --git a/runtime/hsa-runtime/core/inc/intercept_queue.h b/runtime/hsa-runtime/core/inc/intercept_queue.h new file mode 100644 index 0000000000..2ad2abc23b --- /dev/null +++ b/runtime/hsa-runtime/core/inc/intercept_queue.h @@ -0,0 +1,413 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// 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_INTERCEPT_QUEUE_H_ +#define HSA_RUNTIME_CORE_INC_INTERCEPT_QUEUE_H_ + +#include +#include +#include + +#include "core/inc/runtime.h" +#include "core/inc/queue.h" +#include "core/inc/signal.h" +#include "core/inc/interrupt_signal.h" +#include "core/inc/exceptions.h" +#include "core/util/locks.h" + +namespace core { + +// @brief Generic container to forward Queue interfaces into Queue* member. +// Class only has utility as a base type customized Queue wrappers. +class QueueWrapper : public Queue { + public: + std::unique_ptr wrapped; + + explicit QueueWrapper(std::unique_ptr queue) : Queue(), wrapped(std::move(queue)) { + if (Queue::Shared::shared_object() == NULL) { + return; + } + + memcpy(&amd_queue_, &wrapped->amd_queue_, sizeof(amd_queue_t)); + wrapped->set_public_handle(wrapped.get(), public_handle_); + } + + ~QueueWrapper() {} + + hsa_status_t Inactivate() override { return wrapped->Inactivate(); } + uint64_t LoadReadIndexAcquire() override { return wrapped->LoadReadIndexAcquire(); } + uint64_t LoadReadIndexRelaxed() override { return wrapped->LoadReadIndexRelaxed(); } + uint64_t LoadWriteIndexRelaxed() override { return wrapped->LoadWriteIndexRelaxed(); } + uint64_t LoadWriteIndexAcquire() override { return wrapped->LoadWriteIndexAcquire(); } + void StoreReadIndexRelaxed(uint64_t value) override { + return wrapped->StoreReadIndexRelaxed(value); + } + void StoreReadIndexRelease(uint64_t value) override { + return wrapped->StoreReadIndexRelease(value); + } + void StoreWriteIndexRelaxed(uint64_t value) override { + return wrapped->StoreWriteIndexRelaxed(value); + } + void StoreWriteIndexRelease(uint64_t value) override { + return wrapped->StoreWriteIndexRelease(value); + } + uint64_t CasWriteIndexAcqRel(uint64_t expected, uint64_t value) override { + return wrapped->CasWriteIndexAcqRel(expected, value); + } + uint64_t CasWriteIndexAcquire(uint64_t expected, uint64_t value) override { + return wrapped->CasWriteIndexAcquire(expected, value); + } + uint64_t CasWriteIndexRelaxed(uint64_t expected, uint64_t value) override { + return wrapped->CasWriteIndexRelaxed(expected, value); + } + uint64_t CasWriteIndexRelease(uint64_t expected, uint64_t value) override { + return wrapped->CasWriteIndexRelease(expected, value); + } + uint64_t AddWriteIndexAcqRel(uint64_t value) override { + return wrapped->AddWriteIndexAcqRel(value); + } + uint64_t AddWriteIndexAcquire(uint64_t value) override { + return wrapped->AddWriteIndexAcquire(value); + } + uint64_t AddWriteIndexRelaxed(uint64_t value) override { + return wrapped->AddWriteIndexRelaxed(value); + } + uint64_t AddWriteIndexRelease(uint64_t value) override { + return wrapped->AddWriteIndexRelease(value); + } + hsa_status_t SetCUMasking(const uint32_t num_cu_mask_count, const uint32_t* cu_mask) override { + return wrapped->SetCUMasking(num_cu_mask_count, cu_mask); + } + void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) override { + wrapped->ExecutePM4(cmd_data, cmd_size_b); + } + + protected: + void do_set_public_handle(hsa_queue_t* handle) { + public_handle_ = handle; + wrapped->set_public_handle(wrapped.get(), handle); + } +}; + +// @brief Generic container for a proxy queue. +// Presents an proxy packet buffer and doorbell signal for an underlying Queue. Write index +// operations act on the proxy buffer while all other operations pass through to the underlying +// queue. +class QueueProxy : public QueueWrapper { + public: + explicit QueueProxy(std::unique_ptr queue) : QueueWrapper(std::move(queue)) {} + + uint64_t LoadReadIndexAcquire() override { + return atomic::Load(&amd_queue_.read_dispatch_id, std::memory_order_acquire); + } + uint64_t LoadReadIndexRelaxed() override { + return atomic::Load(&amd_queue_.read_dispatch_id, std::memory_order_relaxed); + } + void StoreReadIndexRelaxed(uint64_t value) override { assert(false); } + void StoreReadIndexRelease(uint64_t value) override { assert(false); } + + uint64_t LoadWriteIndexRelaxed() override { + return atomic::Load(&amd_queue_.write_dispatch_id, std::memory_order_relaxed); + } + uint64_t LoadWriteIndexAcquire() override { + return atomic::Load(&amd_queue_.write_dispatch_id, std::memory_order_acquire); + } + void StoreWriteIndexRelaxed(uint64_t value) override { + atomic::Store(&amd_queue_.write_dispatch_id, value, std::memory_order_relaxed); + } + void StoreWriteIndexRelease(uint64_t value) override { + atomic::Store(&amd_queue_.write_dispatch_id, value, std::memory_order_release); + } + uint64_t CasWriteIndexAcqRel(uint64_t expected, uint64_t value) override { + return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_acq_rel); + } + uint64_t CasWriteIndexAcquire(uint64_t expected, uint64_t value) override { + return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_acquire); + } + uint64_t CasWriteIndexRelaxed(uint64_t expected, uint64_t value) override { + return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_relaxed); + } + uint64_t CasWriteIndexRelease(uint64_t expected, uint64_t value) override { + return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected, std::memory_order_release); + } + uint64_t AddWriteIndexAcqRel(uint64_t value) override { + return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_acq_rel); + } + uint64_t AddWriteIndexAcquire(uint64_t value) override { + return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_acquire); + } + uint64_t AddWriteIndexRelaxed(uint64_t value) override { + return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_relaxed); + } + uint64_t AddWriteIndexRelease(uint64_t value) override { + return atomic::Add(&amd_queue_.write_dispatch_id, value, std::memory_order_release); + } +}; + +// @brief Provides packet intercept and rewrite capability for a queue. +// Host-side dispatches are processed during doorbell ring. +// Device-side dispatches are processed as an asynchronous signal event. +class InterceptQueue : public QueueProxy, private LocalSignal, public Signal { + public: + // typedef void(*intercept_packet_writer)(const AqlPacket* pkts, uint64_t pkt_count); + // typedef void(*intercept_handler)(const AqlPacket* pkts, uint64_t pkt_count, uint64_t + // user_pkt_index, void* data, intercept_packet_writer writer); + + std::vector> interceptors; + + explicit InterceptQueue(std::unique_ptr queue); + ~InterceptQueue(); + + void AddInterceptor(hsa_amd_queue_intercept_handler interceptor, void* data) { + interceptors.push_back(std::make_pair(interceptor, data)); + } + + hsa_status_t Inactivate() override { + active_ = false; + return wrapped->Inactivate(); + } + + private: + // Serialize packet interception processing. + KernelMutex lock_; + + // Largest processed packet index. + uint64_t next_packet_; + + // Post interception packet overflow buffer + std::vector overflow_; + + // Index at which async intercept processing was scheduled. + uint64_t retry_index_; + + // Event signal to use for async packet processing and control flag. + InterruptSignal* async_doorbell_; + std::atomic quit_; + + // Indicates queue active/inactive state. + std::atomic active_; + + static const hsa_signal_value_t DOORBELL_MAX = 0xFFFFFFFFFFFFFFFFull; + + static bool HandleAsyncDoorbell(hsa_signal_value_t value, void* arg); + static void PacketWriter(const void* pkts, uint64_t pkt_count); + + bool Submit(const AqlPacket* packets, uint64_t count); + static void Submit(const void* pkts, uint64_t pkt_count, uint64_t user_pkt_index, void* data, + hsa_amd_queue_intercept_packet_writer writer); + + /* + * Remaining Queue and Signal interface definitions. + */ + public: + /// @brief Update signal value using Relaxed semantics + /// + /// @param value Value of signal to update with + void StoreRelaxed(hsa_signal_value_t value); + + /// @brief Update signal value using Release semantics + /// + /// @param value Value of signal to update with + void StoreRelease(hsa_signal_value_t value) { + std::atomic_thread_fence(std::memory_order_release); + StoreRelaxed(value); + } + + /// @brief This operation is illegal + hsa_signal_value_t LoadRelaxed() { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t LoadAcquire() { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t WaitRelaxed(hsa_signal_condition_t condition, hsa_signal_value_t compare_value, + uint64_t timeout, hsa_wait_state_t wait_hint) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t WaitAcquire(hsa_signal_condition_t condition, hsa_signal_value_t compare_value, + uint64_t timeout, hsa_wait_state_t wait_hint) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + void AndRelaxed(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void AndAcquire(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void AndRelease(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void AndAcqRel(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void OrRelaxed(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void OrAcquire(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void OrRelease(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void OrAcqRel(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void XorRelaxed(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void XorAcquire(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void XorRelease(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void XorAcqRel(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void AddRelaxed(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void AddAcquire(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void AddRelease(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void AddAcqRel(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void SubRelaxed(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void SubAcquire(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void SubRelease(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + void SubAcqRel(hsa_signal_value_t value) { assert(false); } + + /// @brief This operation is illegal + hsa_signal_value_t ExchRelaxed(hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t ExchAcquire(hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t ExchRelease(hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t ExchAcqRel(hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t CasRelaxed(hsa_signal_value_t expected, hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t CasAcquire(hsa_signal_value_t expected, hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t CasRelease(hsa_signal_value_t expected, hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t CasAcqRel(hsa_signal_value_t expected, hsa_signal_value_t value) { + assert(false); + return 0; + } + + /// @brief This operation is illegal + hsa_signal_value_t* ValueLocation() const { + assert(false); + return NULL; + } + + /// @brief This operation is illegal + HsaEvent* EopEvent() { + assert(false); + return NULL; + } + + static __forceinline bool IsType(core::Signal* signal) { return signal->IsType(&rtti_id_); } + static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); } + + protected: + bool _IsA(Queue::rtti_t id) const { return id == &rtti_id_; } + + private: + static int rtti_id_; +}; + +} // namespace core + +#endif // HSA_RUNTIME_CORE_INC_INTERCEPT_QUEUE_H_ diff --git a/runtime/hsa-runtime/core/inc/queue.h b/runtime/hsa-runtime/core/inc/queue.h index 5804449faf..1394fd50af 100644 --- a/runtime/hsa-runtime/core/inc/queue.h +++ b/runtime/hsa-runtime/core/inc/queue.h @@ -306,6 +306,10 @@ class Queue : public Checked<0xFA3906A679F9DB49>, hsa_queue_t* public_handle() const { return public_handle_; } + typedef void* rtti_t; + + bool IsType(rtti_t id) { return _IsA(id); } + protected: static void set_public_handle(Queue* ptr, hsa_queue_t* handle) { ptr->do_set_public_handle(handle); @@ -315,6 +319,8 @@ class Queue : public Checked<0xFA3906A679F9DB49>, } hsa_queue_t* public_handle_; + virtual bool _IsA(rtti_t id) const = 0; + private: DISALLOW_COPY_AND_ASSIGN(Queue); }; diff --git a/runtime/hsa-runtime/core/inc/runtime.h b/runtime/hsa-runtime/core/inc/runtime.h index f5bbb5ac55..ac0d4e8d31 100644 --- a/runtime/hsa-runtime/core/inc/runtime.h +++ b/runtime/hsa-runtime/core/inc/runtime.h @@ -50,6 +50,7 @@ #include "core/inc/hsa_ext_interface.h" #include "core/inc/hsa_internal.h" +#include "core/inc/hsa_ext_amd_impl.h" #include "core/inc/agent.h" #include "core/inc/memory_region.h" diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index 0298266a5a..00a4b66b53 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -964,10 +964,8 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // Attempt to trim the maximum number of concurrent waves to allow scratch to fit. // This is somewhat dangerous as it limits the number of concurrent waves from future dispatches // on the queue if those waves use even small amounts of scratch. -#ifndef NDEBUG if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) - fprintf(stderr, "Failed to map requested scratch - reducing queue occupancy.\n"); -#endif + debug_print("Failed to map requested scratch - reducing queue occupancy.\n"); uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU; uint64_t size_per_wave = AlignUp(scratch.size_per_thread * properties_.WaveFrontSize, 1024); uint64_t total_waves = scratch.size / size_per_wave; @@ -994,10 +992,8 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // Failed to allocate minimal scratch assert(scratch.queue_base == NULL && "bad scratch data"); -#ifndef NDEBUG if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) - fprintf(stderr, "Could not allocate scratch for one wave per CU.\n"); -#endif + debug_print("Could not allocate scratch for one wave per CU.\n"); } void GpuAgent::ReleaseQueueScratch(void* base) { diff --git a/runtime/hsa-runtime/core/runtime/host_queue.cpp b/runtime/hsa-runtime/core/runtime/host_queue.cpp index 31ba080171..ca6d2aca3c 100644 --- a/runtime/hsa-runtime/core/runtime/host_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/host_queue.cpp @@ -47,6 +47,7 @@ namespace core { +int HostQueue::rtti_id_ = 0; std::atomic HostQueue::queue_count_(0x80000000); HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, diff --git a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp index f0e5f1278c..fc4238f1a7 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp @@ -47,6 +47,17 @@ #include +// Tools only APIs. +namespace AMD { +hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue, + hsa_amd_queue_intercept_handler callback, + void* user_data); +hsa_status_t hsa_amd_queue_intercept_create( + hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type, + 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); +} + namespace core { HsaApiTable hsa_api_table_; @@ -382,6 +393,8 @@ void HsaApiTable::UpdateAmdExts() { 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; amd_ext_api.hsa_amd_register_system_event_handler_fn = AMD::hsa_amd_register_system_event_handler; + amd_ext_api.hsa_amd_queue_intercept_create_fn = AMD::hsa_amd_queue_intercept_create; + amd_ext_api.hsa_amd_queue_intercept_register_fn = AMD::hsa_amd_queue_intercept_register; } 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 99da16607e..3819e4cdb9 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp @@ -44,8 +44,8 @@ #include #include #include - -#include "hsakmt.h" +#include +#include #include "core/inc/runtime.h" #include "core/inc/agent.h" @@ -56,6 +56,7 @@ #include "core/inc/default_signal.h" #include "core/inc/interrupt_signal.h" #include "core/inc/ipc_signal.h" +#include "core/inc/intercept_queue.h" #include "core/inc/exceptions.h" template @@ -134,12 +135,10 @@ hsa_status_t handleException() { // Rethrow exceptions caught from callbacks. // e.rethrow_nested(); // return HSA_STATUS_ERROR; -#ifndef NDEBUG } catch (const std::exception& e) { - fprintf(stderr, "Unhandled exception: %s\n", e.what()); + debug_print("Unhandled exception: %s\n", e.what()); assert(false && "Unhandled exception."); return HSA_STATUS_ERROR; -#endif } catch (...) { assert(false && "Unhandled exception."); abort(); @@ -798,6 +797,45 @@ hsa_status_t hsa_amd_ipc_signal_attach(const hsa_amd_ipc_signal_t* handle, CATCH; } +// For use by tools only - not in library export table. +hsa_status_t hsa_amd_queue_intercept_create( + hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type, + 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; + hsa_queue_t* lower_queue; + hsa_status_t err = HSA::hsa_queue_create(agent_handle, size, type, callback, data, + private_segment_size, group_segment_size, &lower_queue); + if (err != HSA_STATUS_SUCCESS) return err; + std::unique_ptr lowerQueue(core::Queue::Convert(lower_queue)); + + std::unique_ptr upperQueue(new core::InterceptQueue(std::move(lowerQueue))); + + if (!upperQueue->IsSharedObjectAllocationValid()) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + + *queue = core::Queue::Convert(upperQueue.get()); + upperQueue.release(); + + return HSA_STATUS_SUCCESS; + CATCH; +} + +// For use by tools only - not in library export table. +hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue, + hsa_amd_queue_intercept_handler callback, + void* user_data) { + TRY; + IS_OPEN(); + IS_BAD_PTR(callback); + core::Queue* cmd_queue = core::Queue::Convert(queue); + IS_VALID(cmd_queue); + if (!core::InterceptQueue::IsType(cmd_queue)) return HSA_STATUS_ERROR_INVALID_QUEUE; + core::InterceptQueue* iQueue = static_cast(cmd_queue); + iQueue->AddInterceptor(callback, user_data); + return HSA_STATUS_SUCCESS; + CATCH; +} + hsa_status_t hsa_amd_register_system_event_handler( hsa_amd_event_t type, hsa_status_t (*callback)(const void* event_specific_data, void* data), diff --git a/runtime/hsa-runtime/core/runtime/intercept_queue.cpp b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp new file mode 100644 index 0000000000..387d352578 --- /dev/null +++ b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp @@ -0,0 +1,260 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// 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/intercept_queue.h" +#include "core/util/utils.h" + +namespace core { + +struct InterceptFrame { + InterceptQueue* queue; + uint64_t pkt_index; + size_t interceptor_index; +}; + +static thread_local InterceptFrame Cursor = {nullptr, 0, 0}; + +static const uint16_t kInvalidHeader = (HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); + +static const uint16_t kBarrierHeader = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); + +static const hsa_barrier_and_packet_t kBarrierPacket = {kInvalidHeader, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + +int InterceptQueue::rtti_id_ = 0; + +InterceptQueue::InterceptQueue(std::unique_ptr queue) + : QueueProxy(std::move(queue)), + LocalSignal(0), + Signal(signal()), + next_packet_(0), + retry_index_(0), + quit_(false), + active_(true) { + if (Queue::Shared::shared_object() == NULL) { + return; + // TODO skeely: cleanup queue constructor failure. + // throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "Failed to allocate Queue ABI + // block.\n"); + } + + void* buffer = Runtime::runtime_singleton_->system_allocator()( + wrapped->amd_queue_.hsa_queue.size * sizeof(hsa_agent_dispatch_packet_t), 4096, + MemoryRegion::AllocateNoFlags); + if (buffer == nullptr) throw std::bad_alloc(); + MAKE_NAMED_SCOPE_GUARD(buffGuard, + [&]() { Runtime::runtime_singleton_->system_deallocator()(buffer); }); + amd_queue_.hsa_queue.base_address = buffer; + + // Match the queue's signal ABI block to async_doorbell_'s + // This allows us to use the queue's signal ABI block from devices to trigger async_doorbell while + // host side use jumps directly to the queue's signal implementation. + async_doorbell_ = new InterruptSignal(DOORBELL_MAX); + MAKE_NAMED_SCOPE_GUARD(sigGuard, [&]() { async_doorbell_->DestroySignal(); }); + this->signal_ = async_doorbell_->signal_; + amd_queue_.hsa_queue.doorbell_signal = Signal::Convert(this); + + // Install an async handler for device side dispatches. + auto err = Runtime::runtime_singleton_->SetAsyncSignalHandler( + core::Signal::Convert(async_doorbell_), HSA_SIGNAL_CONDITION_NE, + async_doorbell_->LoadRelaxed(), HandleAsyncDoorbell, this); + if (err != HSA_STATUS_SUCCESS) + throw AMD::hsa_exception(err, "Doorbell handler registration failed.\n"); + + // Install copy submission interceptor. + AddInterceptor(Submit, this); + + sigGuard.Dismiss(); + buffGuard.Dismiss(); +} + +InterceptQueue::~InterceptQueue() { + if (Queue::Shared::shared_object() == NULL) { + return; + } + + active_ = false; + + // Kill the async doorbell handler + // Doorbell may not be used during or after queue destroy, however an interrupt may be in flight. + // Ensure doorbell value is not 0, mark for exit, wake handler and wait for termination value. + async_doorbell_->StoreRelaxed(DOORBELL_MAX); + quit_ = true; + hsa_signal_value_t val = async_doorbell_->ExchRelaxed(1); + if (val != 0) + async_doorbell_->WaitRelaxed(HSA_SIGNAL_CONDITION_EQ, 0, -1, HSA_WAIT_STATE_BLOCKED); + async_doorbell_->DestroySignal(); + + // Release buffer resources + Runtime::runtime_singleton_->system_deallocator()(amd_queue_.hsa_queue.base_address); +} + +bool InterceptQueue::HandleAsyncDoorbell(hsa_signal_value_t value, void* arg) { + InterceptQueue* queue = reinterpret_cast(arg); + if (queue->quit_) { + queue->async_doorbell_->StoreRelaxed(0); + return false; + } + queue->async_doorbell_->StoreRelaxed(DOORBELL_MAX); + queue->StoreRelease(value); + return true; +} + +void InterceptQueue::PacketWriter(const void* pkts, uint64_t pkt_count) { + Cursor.interceptor_index--; + auto& handler = Cursor.queue->interceptors[Cursor.interceptor_index]; + handler.first(pkts, pkt_count, Cursor.pkt_index, handler.second, PacketWriter); +} + +void InterceptQueue::Submit(const void* pkts, uint64_t pkt_count, uint64_t user_pkt_index, + void* data, hsa_amd_queue_intercept_packet_writer writer) { + InterceptQueue* queue = reinterpret_cast(data); + const AqlPacket* packets = (const AqlPacket*)pkts; + + // Submit final packet transform to hardware. + if (queue->Submit(packets, pkt_count)) return; + + // Could not submit final packets, stash for later. + assert(queue->overflow_.empty() && "Packet intercept error: overflow buffer not empty.\n"); + for (uint64_t i = 0; i < pkt_count; i++) queue->overflow_.push_back(packets[i]); +} + +bool InterceptQueue::Submit(const AqlPacket* packets, uint64_t count) { + if (count == 0) return true; + + AqlPacket* ring = reinterpret_cast(wrapped->amd_queue_.hsa_queue.base_address); + uint64_t mask = wrapped->amd_queue_.hsa_queue.size - 1; + + while (true) { + uint64_t write = wrapped->LoadWriteIndexRelaxed(); + uint64_t read = wrapped->LoadReadIndexRelaxed(); + uint64_t free_slots = wrapped->amd_queue_.hsa_queue.size - (write - read); + + // If out of space defer packet insertion. + if (free_slots <= count) { + // If there is not already a pending retry point add one. + if (retry_index_ <= read) { + // Reserve and wait for one slot. + write = wrapped->AddWriteIndexRelaxed(1); + read = write - wrapped->amd_queue_.hsa_queue.size + 1; + while (wrapped->LoadReadIndexRelaxed() < read) os::YieldThread(); + + // Submit barrer which will wake async queue processing. + ring[write & mask].barrier_and = kBarrierPacket; + ring[write & mask].barrier_and.completion_signal = Signal::Convert(async_doorbell_); + atomic::Store(&ring[write & mask].barrier_and.header, kBarrierHeader, + std::memory_order_release); + HSA::hsa_signal_store_screlease(wrapped->amd_queue_.hsa_queue.doorbell_signal, write); + + // Record the retry point + retry_index_ = write; + } + return false; + } + + // Attempt to reserve useable queue space + uint64_t new_write = wrapped->CasWriteIndexRelaxed(write, write + count); + if (new_write == write) { + AqlPacket first = packets[0]; + uint16_t header = first.dispatch.header; + first.dispatch.header = kInvalidHeader; + + ring[write & mask] = first; + for (uint64_t i = 1; i < count; i++) ring[(write + i) & mask] = packets[i]; + atomic::Store(&ring[write & mask].dispatch.header, header, std::memory_order_release); + HSA::hsa_signal_store_screlease(wrapped->amd_queue_.hsa_queue.doorbell_signal, + write + count - 1); + + return true; + } + } +} + +void InterceptQueue::StoreRelaxed(hsa_signal_value_t value) { + if (!active_) return; + + // If called recursively defer to async doorbell thread. + if (Cursor.queue != nullptr) { + debug_print("Likely incorrect queue use observed in an interceptor.\n"); + async_doorbell_->StoreRelaxed(value); + return; + } + + ScopedAcquire lock(&lock_); + + // Submit overflow packets. + if (!overflow_.empty()) { + if (!Submit(&overflow_[0], overflow_.size())) return; + overflow_.clear(); + } + + Cursor.queue = this; + + AqlPacket* ring = reinterpret_cast(amd_queue_.hsa_queue.base_address); + uint64_t mask = wrapped->amd_queue_.hsa_queue.size - 1; + + // Loop over valid packets and process. + uint64_t end = LoadWriteIndexAcquire(); + uint64_t i; + for (i = next_packet_; i < end; i++) { + if (!ring[i & mask].IsValid()) break; + + // Process callbacks. + Cursor.interceptor_index = interceptors.size() - 1; + Cursor.pkt_index = i; + auto& handler = interceptors[Cursor.interceptor_index]; + handler.first(&ring[i & mask], 1, i, handler.second, PacketWriter); + + // Invalidate consumed packet + atomic::Store(&ring[i & mask].dispatch.header, kInvalidHeader, std::memory_order_release); + } + + next_packet_ = i; + Cursor.queue = nullptr; +} + +} // namespace core diff --git a/runtime/hsa-runtime/core/runtime/runtime.cpp b/runtime/hsa-runtime/core/runtime/runtime.cpp index 59bb32be23..c02c3a635c 100644 --- a/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -1320,12 +1320,10 @@ void Runtime::LoadTools() { add = (tool_add_t)os::GetExportAddress(tool, "AddAgent"); if (add) add(this); } -#ifndef NDEBUG else { if (flag().report_tool_load_failures()) - fprintf(stderr, "Tool lib \"%s\" failed to load.\n", names[i].c_str()); + debug_print("Tool lib \"%s\" failed to load.\n", names[i].c_str()); } -#endif } } } diff --git a/runtime/hsa-runtime/core/util/utils.h b/runtime/hsa-runtime/core/util/utils.h index 0e9a2aa8bd..2da2d09668 100755 --- a/runtime/hsa-runtime/core/util/utils.h +++ b/runtime/hsa-runtime/core/util/utils.h @@ -78,12 +78,15 @@ static __forceinline void _aligned_free(void* ptr) { return free(ptr); } #elif defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) #include "intrin.h" #define __ALIGNED__(x) __declspec(align(x)) -#if (_MSC_VER < 1800) +#if (_MSC_VER < 1800) // < VS 2013 static __forceinline unsigned long long int strtoull(const char* str, char** endptr, int base) { return static_cast(_strtoui64(str, endptr, base)); } #endif +#if (_MSC_VER < 1900) // < VS 2015 +#define thread_local __declspec(thread) +#endif #else #error "Compiler and/or processor not identified." #endif @@ -102,7 +105,18 @@ static __forceinline unsigned long long int strtoull(const char* str, if (!(exp)) \ fprintf(stderr, "Warning: " STRING(exp) " in %s, " __FILE__ ":" STRING(__LINE__) "\n", \ __PRETTY_FUNCTION__); \ - } while (false); + } while (false) +#endif + +#ifdef NDEBUG +#define debug_print(fmt, ...) \ + do { \ + } while (false) +#else +#define debug_print(fmt, ...) \ + do { \ + fprintf(stderr, fmt, ##__VA_ARGS__); \ + } while (false) #endif // A macro to disallow the copy and move constructor and operator= functions diff --git a/runtime/hsa-runtime/inc/hsa_api_trace.h b/runtime/hsa-runtime/inc/hsa_api_trace.h index 455e3399ec..fcdfc1e5dc 100644 --- a/runtime/hsa-runtime/inc/hsa_api_trace.h +++ b/runtime/hsa-runtime/inc/hsa_api_trace.h @@ -81,6 +81,19 @@ static inline uint32_t Min(const uint32_t a, const uint32_t b) { return (a > b) ? b : a; } +// Declarations of APIs intended for use only by tools. +typedef void (*hsa_amd_queue_intercept_packet_writer)(const void* pkts, uint64_t pkt_count); +typedef void (*hsa_amd_queue_intercept_handler)(const void* pkts, uint64_t pkt_count, + uint64_t user_pkt_index, void* data, + hsa_amd_queue_intercept_packet_writer writer); +hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue, + hsa_amd_queue_intercept_handler callback, + void* user_data); +hsa_status_t hsa_amd_queue_intercept_create( + hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type, + 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); + // Structure of Version used to identify an instance of Api table // Must be the first member (offsetof == 0) of all API tables. // This is the root of the table passing ABI. @@ -170,6 +183,8 @@ struct AmdExtTable { decltype(hsa_amd_ipc_signal_create)* hsa_amd_ipc_signal_create_fn; decltype(hsa_amd_ipc_signal_attach)* hsa_amd_ipc_signal_attach_fn; decltype(hsa_amd_register_system_event_handler)* hsa_amd_register_system_event_handler_fn; + decltype(hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn; + decltype(hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn; }; // Table to export HSA Core Runtime Apis