From 6c556002d84ece04edb1d4875d5562c609db92d6 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Fri, 14 Feb 2020 21:43:32 -0600 Subject: [PATCH] Correct scratch retry logic. scratch_used_large_ was uninitialized leading to the observed hang. DynamicScratchHandler would wait for a large scratch release despite no large scratch having yet been allocated. Fixes . The patch also removes a potential race between AddScratchNotifier and ReleaseQueueScratch. The race condition does not exist today since both scratch alloc and release run on the same thread. The changes will prevent this potential race from manifesting if the async event handler is ever updated to use multiple threads. Also enhances scratch occupancy reduction reporting. Reporting now prints the initial request size as well as the allocated size and the effect on occupancy this has. Occupancy is computed in terms of the requesting dispatch grid size so may be >100%. Change-Id: I0fc5ee01467ff4c29bdd25d545177c97862c3bd9 --- runtime/hsa-runtime/core/inc/amd_gpu_agent.h | 28 ++++++++--------- .../core/runtime/amd_aql_queue.cpp | 22 ++++++++----- .../core/runtime/amd_gpu_agent.cpp | 31 +++++++++++-------- 3 files changed, 45 insertions(+), 36 deletions(-) diff --git a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index a538518afc..59a5ce840b 100644 --- a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -71,6 +71,8 @@ struct ScratchInfo { ptrdiff_t queue_process_offset; bool large; bool retry; + hsa_signal_t queue_retry; + uint64_t wanted_slots; }; // @brief Interface to represent a GPU agent. @@ -265,22 +267,8 @@ class GpuAgent : public GpuAgentInt { // @brief Override from amd::GpuAgentInt. void ReleaseQueueScratch(ScratchInfo& scratch) override; - // @brief Register signal for notification when scratch may become available. - // @p signal is notified by OR'ing with @p value. - void AddScratchNotifier(hsa_signal_t signal, hsa_signal_value_t value) { - ScopedAcquire lock(&scratch_lock_); - scratch_notifiers_[signal] = value; - } - - // @brief Deregister scratch notification signal. - void RemoveScratchNotifier(hsa_signal_t signal) { - ScopedAcquire lock(&scratch_lock_); - scratch_notifiers_.erase(signal); - } - // @brief Override from amd::GpuAgentInt. - void TranslateTime(core::Signal* signal, - hsa_amd_profiling_dispatch_time_t& time) override; + void TranslateTime(core::Signal* signal, hsa_amd_profiling_dispatch_time_t& time) override; // @brief Override from amd::GpuAgentInt. void TranslateTime(core::Signal* signal, hsa_amd_profiling_async_copy_time_t& time) override; @@ -494,6 +482,16 @@ class GpuAgent : public GpuAgentInt { // @brief Setup GWS accessing queue. void InitGWS(); + // @brief Register signal for notification when scratch may become available. + // @p signal is notified by OR'ing with @p value. + bool AddScratchNotifier(hsa_signal_t signal, hsa_signal_value_t value) { + if (signal.handle != 0) return false; + scratch_notifiers_[signal] = value; + return true; + } + + // @brief Deregister scratch notification signals. + void ClearScratchNotifiers() { scratch_notifiers_.clear(); } // Bind index of peer device that is connected via xGMI links lazy_ptr& GetXgmiBlit(const core::Agent& peer_agent); diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index b012f67522..4ce4f9093b 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -189,9 +189,6 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr 0); #endif - // Initialize scratch memory related entities - InitScratchSRD(); - // Set group and private memory apertures in amd_queue_. auto& regions = agent->regions(); @@ -257,6 +254,11 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr assert(Signal != nullptr && "Should have thrown!\n"); amd_queue_.queue_inactive_signal = core::DefaultSignal::Convert(Signal); } + + // Initialize scratch memory related entities + queue_scratch_.queue_retry = amd_queue_.queue_inactive_signal; + InitScratchSRD(); + if (AMD::hsa_amd_signal_async_handler(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE, 0, DynamicScratchHandler, this) != HSA_STATUS_SUCCESS) throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, @@ -729,7 +731,6 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { if ((queue->dynamicScratchState & ERROR_HANDLER_SCRATCH_RETRY) == ERROR_HANDLER_SCRATCH_RETRY) { queue->dynamicScratchState &= ~ERROR_HANDLER_SCRATCH_RETRY; - queue->agent_->RemoveScratchNotifier(queue->amd_queue_.queue_inactive_signal); changeWait = true; waitVal = 0; HSA::hsa_signal_and_relaxed(queue->amd_queue_.queue_inactive_signal, ~0x8000000000000000ull); @@ -771,18 +772,23 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { uint32_t scratch_request = pkt.dispatch.private_segment_size; + const uint32_t MaxScratchSlots = + (queue->amd_queue_.max_cu_id + 1) * queue->agent_->properties().MaxSlotsScratchCU; + scratch.size_per_thread = scratch_request; scratch.lanes_per_wave = (error_code & 0x400) ? 32 : 64; // Align whole waves to 1KB. scratch.size_per_thread = AlignUp(scratch.size_per_thread, 1024 / scratch.lanes_per_wave); - scratch.size = scratch.size_per_thread * (queue->amd_queue_.max_cu_id + 1) * - queue->agent_->properties().MaxSlotsScratchCU * scratch.lanes_per_wave; + scratch.size = scratch.size_per_thread * MaxScratchSlots * scratch.lanes_per_wave; +#ifndef NDEBUG + scratch.wanted_slots = ((uint64_t(pkt.dispatch.grid_size_x) * pkt.dispatch.grid_size_y) * + pkt.dispatch.grid_size_z) / scratch.lanes_per_wave; + scratch.wanted_slots = Min(scratch.wanted_slots, MaxScratchSlots); +#endif queue->agent_->AcquireQueueScratch(scratch); if (scratch.retry) { - queue->agent_->AddScratchNotifier(queue->amd_queue_.queue_inactive_signal, - 0x8000000000000000ull); queue->dynamicScratchState |= ERROR_HANDLER_SCRATCH_RETRY; changeWait = true; waitVal = error_code; diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index b2acd54afe..d0bd897043 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -77,6 +77,7 @@ GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props) : GpuAgentInt(node), properties_(node_props), current_coherency_type_(HSA_AMD_COHERENCY_TYPE_COHERENT), + scratch_used_large_(0), queues_(), local_region_(NULL), is_kv_device_(false), @@ -936,7 +937,7 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, } // Allocate scratch memory - ScratchInfo scratch; + ScratchInfo scratch = {0}; if (private_segment_size == UINT_MAX) { private_segment_size = (profile_ == HSA_PROFILE_BASE) ? 0 : scratch_per_thread_; } @@ -1038,15 +1039,16 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // Retry if large may yield needed space. if (scratch_used_large_ != 0) { - scratch.retry = true; + if (AddScratchNotifier(scratch.queue_retry, 0x8000000000000000ull)) scratch.retry = true; return; } // Attempt to trim the maximum number of concurrent waves to allow scratch to fit. if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) - debug_print("Failed to map requested scratch - reducing queue occupancy.\n"); - uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU; - uint64_t total_waves = scratch.size / size_per_wave; + debug_print("Failed to map requested scratch (%ld) - reducing queue occupancy.\n", + scratch.size); + const uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU; + const uint64_t total_waves = scratch.size / size_per_wave; uint64_t waves_per_cu = total_waves / num_cus; while (waves_per_cu != 0) { size_t size = waves_per_cu * num_cus * size_per_wave; @@ -1058,12 +1060,14 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // Scratch allocated and either full profile or map succeeded. scratch.queue_base = base; scratch.size = size; - scratch.queue_process_offset = - (need_queue_scratch_base) - ? uintptr_t(scratch.queue_base) - : uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base()); + scratch.queue_process_offset = (need_queue_scratch_base) + ? uintptr_t(scratch.queue_base) + : uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base()); scratch.large = true; scratch_used_large_ += scratch.size; + if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) + debug_print(" %ld scratch mapped, %.2f%% occupancy.\n", scratch.size, + float(waves_per_cu * num_cus) / scratch.wanted_slots * 100.0f); return; } scratch_pool_.free(base); @@ -1073,7 +1077,7 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // Failed to allocate minimal scratch assert(scratch.queue_base == nullptr && "bad scratch data"); if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) - debug_print("Could not allocate scratch for one wave per CU.\n"); + debug_print(" Could not allocate scratch for one wave per CU.\n"); } void GpuAgent::ReleaseQueueScratch(ScratchInfo& scratch) { @@ -1093,12 +1097,13 @@ void GpuAgent::ReleaseQueueScratch(ScratchInfo& scratch) { if (scratch.large) scratch_used_large_ -= scratch.size; // Notify waiters that additional scratch may be available. - for (auto notifier : scratch_notifiers_) + for (auto notifier : scratch_notifiers_) { HSA::hsa_signal_or_relaxed(notifier.first, notifier.second); + } + ClearScratchNotifiers(); } -void GpuAgent::TranslateTime(core::Signal* signal, - hsa_amd_profiling_dispatch_time_t& time) { +void GpuAgent::TranslateTime(core::Signal* signal, hsa_amd_profiling_dispatch_time_t& time) { uint64_t start, end; signal->GetRawTs(false, start, end); // Order is important, we want to translate the end time first to ensure that packet duration is