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