From fa317f8c41545055e86a21b1966488bbc8ce2697 Mon Sep 17 00:00:00 2001 From: David Yat Sin Date: Wed, 15 Nov 2023 16:34:19 +0000 Subject: [PATCH] Re-arrange and rename scratch elements that are used with main scratch Change-Id: I4c1ff8cf4121a06b586fe49c70400226506bf95e --- runtime/hsa-runtime/core/inc/amd_gpu_agent.h | 16 ++- runtime/hsa-runtime/core/inc/queue.h | 14 +++ runtime/hsa-runtime/core/inc/scratch_cache.h | 76 +++++------ .../core/runtime/amd_aql_queue.cpp | 77 ++++++------ .../core/runtime/amd_gpu_agent.cpp | 119 +++++++++--------- 5 files changed, 163 insertions(+), 139 deletions(-) diff --git a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index 39ed8c04a6..b39e449f85 100644 --- a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -98,17 +98,17 @@ class GpuAgentInt : public core::Agent { void* data), void* data) const = 0; - // @brief Carve scratch memory from scratch pool. + // @brief Carve scratch memory for main from scratch pool. // // @param [in/out] scratch Structure to be populated with the carved memory // information. - virtual void AcquireQueueScratch(ScratchInfo& scratch) = 0; + virtual void AcquireQueueMainScratch(ScratchInfo& scratch) = 0; - // @brief Release scratch memory back to scratch pool. + // @brief Release scratch memory from main back to scratch pool. // // @param [in/out] scratch Scratch memory previously acquired with call to - // ::AcquireQueueScratch. - virtual void ReleaseQueueScratch(ScratchInfo& base) = 0; + // ::AcquireQueueMainScratch. + virtual void ReleaseQueueMainScratch(ScratchInfo& base) = 0; // @brief Translate the kernel start and end dispatch timestamp from agent // domain to host domain. @@ -267,10 +267,8 @@ class GpuAgent : public GpuAgentInt { void GWSRelease(); // @brief Override from AMD::GpuAgentInt. - void AcquireQueueScratch(ScratchInfo& scratch) override; - - // @brief Override from AMD::GpuAgentInt. - void ReleaseQueueScratch(ScratchInfo& scratch) override; + void AcquireQueueMainScratch(ScratchInfo& scratch) override; + void ReleaseQueueMainScratch(ScratchInfo& scratch) override; // @brief Override from AMD::GpuAgentInt. void TranslateTime(core::Signal* signal, hsa_amd_profiling_dispatch_time_t& time) override; diff --git a/runtime/hsa-runtime/core/inc/queue.h b/runtime/hsa-runtime/core/inc/queue.h index e8ac0484e4..468c073282 100644 --- a/runtime/hsa-runtime/core/inc/queue.h +++ b/runtime/hsa-runtime/core/inc/queue.h @@ -90,6 +90,20 @@ struct AqlPacket { (type(header) != HSA_PACKET_TYPE_INVALID)); } + void __forceinline AssertIsDispatchAndNeedsScratch() const { + assert(IsValid(packet.header) && "Invalid packet in dynamic scratch handler."); + assert(type(packet.header) == HSA_PACKET_TYPE_KERNEL_DISPATCH && + "Invalid packet in dynamic scratch handler."); + + assert((dispatch.workgroup_size_x != 0) && (dispatch.workgroup_size_y != 0) && + (dispatch.workgroup_size_z != 0) && "Invalid dispatch dimension."); + + assert((dispatch.private_segment_size != 0) && + "Scratch memory request from packet with no scratch demand. Possible bad kernel code " + "object."); + return; + } + std::string string() const { std::stringstream string; uint8_t t = type(packet.header); diff --git a/runtime/hsa-runtime/core/inc/scratch_cache.h b/runtime/hsa-runtime/core/inc/scratch_cache.h index 792e273516..69e9df8f2f 100644 --- a/runtime/hsa-runtime/core/inc/scratch_cache.h +++ b/runtime/hsa-runtime/core/inc/scratch_cache.h @@ -86,22 +86,26 @@ class ScratchCache { // @brief Contains scratch memory information. struct ScratchInfo { - void* queue_base; - // Size to fill the machine with size_per_thread - size_t size; // Size to satisfy the present dispatch without throttling. size_t dispatch_size; - size_t size_per_thread; - uint32_t lanes_per_wave; - uint32_t waves_per_group; - uint64_t wanted_slots; - uint32_t mem_alignment_size; - bool cooperative; - ptrdiff_t queue_process_offset; + uint64_t dispatch_slots; + bool large; + size_t use_once_limit; + bool async_reclaim; // This version of CP FW supports async_reclaim bool retry; + uint32_t mem_alignment_size; // Populated into SRD + bool cooperative; hsa_signal_t queue_retry; - ScratchCache::ref_t scratch_node; + + // Size to fill the main_scratch with size_per_thread + size_t main_size; + size_t main_size_per_thread; // Populated into SRD + uint32_t main_lanes_per_wave; // Populated into SRD + uint32_t main_waves_per_group; // Used during waves reduction + void* main_queue_base; + ptrdiff_t main_queue_process_offset; + ScratchCache::ref_t main_scratch_node; }; ScratchCache(const ScratchCache& rhs) = delete; @@ -113,17 +117,17 @@ class ScratchCache { ~ScratchCache() { assert(map.empty() && "ScratchCache not empty at shutdown."); } - bool alloc(ScratchInfo& info) { - ref_t it = map.upper_bound(info.size - 1); + bool allocMain(ScratchInfo& info) { + ref_t it = map.upper_bound(info.main_size - 1); if (it == map.end()) return false; // Small requests must have an exact size match and be small. if (!info.large) { - while ((it != map.end()) && (it->first == info.size)) { + while ((it != map.end()) && (it->first == info.main_size)) { if (it->second.isFree() && (!it->second.large)) { it->second.alloc(); - info.queue_base = it->second.base; - info.scratch_node = it; + info.main_queue_base = it->second.base; + info.main_scratch_node = it; available_bytes_ -= it->first; return true; } @@ -136,8 +140,8 @@ class ScratchCache { while (it != map.end()) { if (it->second.isFree()) { it->second.alloc(); - info.queue_base = it->second.base; - info.scratch_node = it; + info.main_queue_base = it->second.base; + info.main_scratch_node = it; available_bytes_ -= it->first; return true; } @@ -146,8 +150,8 @@ class ScratchCache { return false; } - void free(ScratchInfo& info) { - if (info.scratch_node == map.end()) { + void freeMain(ScratchInfo& info) { + if (info.main_scratch_node == map.end()) { // This is reserved scratch memory. Do not de-allocate, just mark it as free. assert(!reserved_.second.isFree() && "free called when reserved node already free."); reserved_.second.free(); @@ -155,8 +159,8 @@ class ScratchCache { return; } - assert(!info.scratch_node->second.isFree() && "free called on free scratch node."); - auto it = info.scratch_node; + assert(!info.main_scratch_node->second.isFree() && "free called on free scratch node."); + auto it = info.main_scratch_node; if (it->second.trimPending()) { dealloc(it->second.base, it->first, it->second.large); map.erase(it); @@ -166,6 +170,16 @@ class ScratchCache { available_bytes_ += it->first; } + void insertMain(ScratchInfo& info) { + node n; + n.base = info.main_queue_base; + n.large = info.large; + n.alloc(); + + auto it = map.insert(std::make_pair(info.main_size, n)); + info.main_scratch_node = it; + } + bool trim(bool trim_nodes_in_use) { bool ret = !map.empty(); auto it = map.begin(); @@ -184,16 +198,6 @@ class ScratchCache { return ret; } - void insert(ScratchInfo& info) { - node n; - n.base = info.queue_base; - n.large = info.large; - n.alloc(); - - auto it = map.insert(std::make_pair(info.size, n)); - info.scratch_node = it; - } - size_t free_bytes() const { return available_bytes_; } size_t reserved_bytes() const { return reserved_.first; } @@ -210,16 +214,16 @@ class ScratchCache { } bool use_reserved(ScratchInfo& info) { - if (!reserved_.second.isFree() || info.size > reserved_.first) { + if (!reserved_.second.isFree() || info.main_size > reserved_.first) { debug_print("reserved node is already in use or too small (requested:%ld reserved:%ld)\n", - info.size, reserved_.first); + info.main_size, reserved_.first); return false; } reserved_.second.large = info.large; reserved_.second.alloc(); - info.queue_base = reserved_.second.base; + info.main_queue_base = reserved_.second.base; // Special case to indicate that this node is reserved memory - info.scratch_node = map.end(); + info.main_scratch_node = map.end(); available_bytes_ -= reserved_.first; return true; } diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 01b18552f1..1a2c56f724 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -348,7 +348,7 @@ AqlQueue::~AqlQueue() { } Inactivate(); - agent_->ReleaseQueueScratch(queue_scratch_); + agent_->ReleaseQueueMainScratch(queue_scratch_); FreeRegisteredRingBuffer(); exception_signal_->DestroySignal(); HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal); @@ -798,11 +798,11 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { if ((queue->dynamicScratchState & ERROR_HANDLER_TERMINATE) != ERROR_HANDLER_TERMINATE) { if (error_code == 512) { // Large scratch reclaim auto& scratch = queue->queue_scratch_; - queue->agent_->ReleaseQueueScratch(scratch); - scratch.queue_base = nullptr; - scratch.size = 0; - scratch.size_per_thread = 0; - scratch.queue_process_offset = 0; + queue->agent_->ReleaseQueueMainScratch(scratch); + scratch.main_queue_base = nullptr; + scratch.main_size = 0; + scratch.main_size_per_thread = 0; + scratch.main_queue_process_offset = 0; queue->InitScratchSRD(); HSA::hsa_signal_store_relaxed(queue->amd_queue_.queue_inactive_signal, 0); @@ -819,7 +819,7 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { // Insufficient scratch - recoverable, don't process dynamic scratch if errors are present. auto& scratch = queue->queue_scratch_; - queue->agent_->ReleaseQueueScratch(scratch); + queue->agent_->ReleaseQueueMainScratch(scratch); uint64_t pkt_slot_idx = queue->amd_queue_.read_dispatch_id & (queue->amd_queue_.hsa_queue.size - 1); @@ -847,13 +847,14 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { const uint32_t MaxScratchSlots = AlignUp(cu_count, engines) * queue->agent_->properties().MaxSlotsScratchCU; - scratch.size_per_thread = scratch_request; - scratch.lanes_per_wave = (error_code & 0x400) ? 32 : 64; + scratch.main_size_per_thread = scratch_request; + scratch.main_lanes_per_wave = (error_code & 0x400) ? 32 : 64; - scratch.size_per_thread = - AlignUp(scratch.size_per_thread, scratch.mem_alignment_size / scratch.lanes_per_wave); + scratch.main_size_per_thread = AlignUp( + scratch.main_size_per_thread, scratch.mem_alignment_size / scratch.main_lanes_per_wave); - scratch.size = scratch.size_per_thread * MaxScratchSlots * scratch.lanes_per_wave; + scratch.main_size = + scratch.main_size_per_thread * MaxScratchSlots * scratch.main_lanes_per_wave; // Smaller dispatches may not need to reach full device occupancy. // For these we need to ensure that the scratch we give doesn't restrict the dispatch even @@ -862,8 +863,8 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { (uint64_t(pkt.dispatch.workgroup_size_x) * pkt.dispatch.workgroup_size_y) * pkt.dispatch.workgroup_size_z; uint64_t waves_per_group = - (lanes_per_group + scratch.lanes_per_wave - 1) / scratch.lanes_per_wave; - scratch.waves_per_group = waves_per_group; + (lanes_per_group + scratch.main_lanes_per_wave - 1) / scratch.main_lanes_per_wave; + scratch.main_waves_per_group = waves_per_group; uint64_t groups = ((uint64_t(pkt.dispatch.grid_size_x) + pkt.dispatch.workgroup_size_x - 1) / pkt.dispatch.workgroup_size_x) * @@ -891,14 +892,14 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { // Populate all engines at max group occupancy, then clip down to device limits. groups = maxGroupsPerEngine * engines; - scratch.wanted_slots = groups * waves_per_group; - scratch.wanted_slots = Min(scratch.wanted_slots, uint64_t(MaxScratchSlots)); + scratch.dispatch_slots = groups * waves_per_group; + scratch.dispatch_slots = Min(scratch.dispatch_slots, uint64_t(MaxScratchSlots)); scratch.dispatch_size = - scratch.size_per_thread * scratch.wanted_slots * scratch.lanes_per_wave; + scratch.main_size_per_thread * scratch.dispatch_slots * scratch.main_lanes_per_wave; scratch.cooperative = (queue->amd_queue_.hsa_queue.type == HSA_QUEUE_TYPE_COOPERATIVE); - queue->agent_->AcquireQueueScratch(scratch); + queue->agent_->AcquireQueueMainScratch(scratch); if (scratch.retry) { queue->dynamicScratchState |= ERROR_HANDLER_SCRATCH_RETRY; @@ -906,7 +907,7 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { waitVal = error_code; } else { // Out of scratch - promote error - if (scratch.queue_base == nullptr) { + if (scratch.main_queue_base == nullptr) { errorCode = HSA_STATUS_ERROR_OUT_OF_RESOURCES; } else { // Mark large scratch allocation for single use. @@ -1294,7 +1295,7 @@ void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) { void AqlQueue::FillBufRsrcWord0() { SQ_BUF_RSRC_WORD0 srd0; - uintptr_t scratch_base = uintptr_t(queue_scratch_.queue_base); + uintptr_t scratch_base = uintptr_t(queue_scratch_.main_queue_base); srd0.bits.BASE_ADDRESS = uint32_t(scratch_base); amd_queue_.scratch_resource_descriptor[0] = srd0.u32All; @@ -1305,7 +1306,7 @@ void AqlQueue::FillBufRsrcWord1() { uint32_t scratch_base_hi = 0; #ifdef HSA_LARGE_MODEL - uintptr_t scratch_base = uintptr_t(queue_scratch_.queue_base); + uintptr_t scratch_base = uintptr_t(queue_scratch_.main_queue_base); scratch_base_hi = uint32_t(scratch_base >> 32); #endif @@ -1322,7 +1323,7 @@ void AqlQueue::FillBufRsrcWord1_Gfx11() { uint32_t scratch_base_hi = 0; #ifdef HSA_LARGE_MODEL - uintptr_t scratch_base = uintptr_t(queue_scratch_.queue_base); + uintptr_t scratch_base = uintptr_t(queue_scratch_.main_queue_base); scratch_base_hi = uint32_t(scratch_base >> 32); #endif @@ -1339,7 +1340,7 @@ void AqlQueue::FillBufRsrcWord2() { const uint32_t num_xcc = agent_props.NumXcc; // report size per XCC - srd2.bits.NUM_RECORDS = uint32_t(queue_scratch_.size / num_xcc); + srd2.bits.NUM_RECORDS = uint32_t(queue_scratch_.main_size / num_xcc); amd_queue_.scratch_resource_descriptor[2] = srd2.u32All; } @@ -1405,7 +1406,7 @@ void AqlQueue::FillBufRsrcWord3_Gfx11() { // Set concurrent wavefront limits only when scratch is being used. void AqlQueue::FillComputeTmpRingSize() { COMPUTE_TMPRING_SIZE tmpring_size = {}; - if (queue_scratch_.size == 0) { + if (queue_scratch_.main_size == 0) { amd_queue_.compute_tmpring_size = tmpring_size.u32All; return; } @@ -1419,13 +1420,14 @@ void AqlQueue::FillComputeTmpRingSize() { // Scratch is allocated program COMPUTE_TMPRING_SIZE register // Scratch Size per Wave is specified in terms of kilobytes - uint32_t wave_scratch = (((queue_scratch_.lanes_per_wave * queue_scratch_.size_per_thread) + - queue_scratch_.mem_alignment_size - 1) / - queue_scratch_.mem_alignment_size); + uint32_t wave_scratch = + (((queue_scratch_.main_lanes_per_wave * queue_scratch_.main_size_per_thread) + + queue_scratch_.mem_alignment_size - 1) / + queue_scratch_.mem_alignment_size); tmpring_size.bits.WAVESIZE = wave_scratch; assert(wave_scratch == tmpring_size.bits.WAVESIZE && "WAVESIZE Overflow."); - uint32_t num_waves = - (queue_scratch_.size / num_xcc) / (tmpring_size.bits.WAVESIZE * queue_scratch_.mem_alignment_size); + uint32_t num_waves = (queue_scratch_.main_size / num_xcc) / + (tmpring_size.bits.WAVESIZE * queue_scratch_.mem_alignment_size); tmpring_size.bits.WAVES = std::min(num_waves, max_scratch_waves); amd_queue_.compute_tmpring_size = tmpring_size.u32All; @@ -1436,7 +1438,7 @@ void AqlQueue::FillComputeTmpRingSize() { // Set concurrent wavefront limits only when scratch is being used. void AqlQueue::FillComputeTmpRingSize_Gfx11() { COMPUTE_TMPRING_SIZE_GFX11 tmpring_size = {}; - if (queue_scratch_.size == 0) { + if (queue_scratch_.main_size == 0) { amd_queue_.compute_tmpring_size = tmpring_size.u32All; return; } @@ -1450,15 +1452,16 @@ void AqlQueue::FillComputeTmpRingSize_Gfx11() { // Scratch is allocated program COMPUTE_TMPRING_SIZE register // Scratch Size per Wave is specified in terms of kilobytes - uint32_t wave_scratch = (((queue_scratch_.lanes_per_wave * queue_scratch_.size_per_thread) + - queue_scratch_.mem_alignment_size - 1) / - queue_scratch_.mem_alignment_size); + uint32_t wave_scratch = + (((queue_scratch_.main_lanes_per_wave * queue_scratch_.main_size_per_thread) + + queue_scratch_.mem_alignment_size - 1) / + queue_scratch_.mem_alignment_size); tmpring_size.bits.WAVESIZE = wave_scratch; assert(wave_scratch == tmpring_size.bits.WAVESIZE && "WAVESIZE Overflow."); uint32_t num_waves = - queue_scratch_.size / (tmpring_size.bits.WAVESIZE * queue_scratch_.mem_alignment_size); + queue_scratch_.main_size / (tmpring_size.bits.WAVESIZE * queue_scratch_.mem_alignment_size); // For GFX11 we specify number of waves per engine instead of total num_waves /= agent_->properties().NumShaderBanks; @@ -1494,18 +1497,18 @@ void AqlQueue::InitScratchSRD() { } // Populate flat scratch parameters in amd_queue_. - amd_queue_.scratch_backing_memory_location = queue_scratch_.queue_process_offset; + amd_queue_.scratch_backing_memory_location = queue_scratch_.main_queue_process_offset; const auto& agent_props = agent_->properties(); const uint32_t num_xcc = agent_props.NumXcc; // report size per XCC - amd_queue_.scratch_backing_memory_byte_size = queue_scratch_.size / num_xcc; + amd_queue_.scratch_backing_memory_byte_size = queue_scratch_.main_size / num_xcc; // For backwards compatibility this field records the per-lane scratch // for a 64 lane wavefront. If scratch was allocated for 32 lane waves // then the effective size for a 64 lane wave is halved. amd_queue_.scratch_wave64_lane_byte_size = - uint32_t((queue_scratch_.size_per_thread * queue_scratch_.lanes_per_wave) / 64); + uint32_t((queue_scratch_.main_size_per_thread * queue_scratch_.main_lanes_per_wave) / 64); return; } diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index b8170a5803..e12cde29c2 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -1555,24 +1555,24 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } - scratch.lanes_per_wave = 64; - scratch.size_per_thread = AlignUp(private_segment_size, 1024 / scratch.lanes_per_wave); - if (scratch.size_per_thread > 262128) { + scratch.main_lanes_per_wave = 64; + scratch.main_size_per_thread = AlignUp(private_segment_size, 1024 / scratch.main_lanes_per_wave); + if (scratch.main_size_per_thread > 262128) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } - scratch.size_per_thread = private_segment_size; + scratch.main_size_per_thread = private_segment_size; const uint32_t num_cu = properties_.NumFComputeCores / properties_.NumSIMDPerCU; - scratch.size = - scratch.size_per_thread * properties_.MaxSlotsScratchCU * scratch.lanes_per_wave * num_cu; - scratch.queue_base = nullptr; - scratch.queue_process_offset = 0; + scratch.main_size = scratch.main_size_per_thread * properties_.MaxSlotsScratchCU * + scratch.main_lanes_per_wave * num_cu; + scratch.main_queue_base = nullptr; + scratch.main_queue_process_offset = 0; - MAKE_NAMED_SCOPE_GUARD(scratchGuard, [&]() { ReleaseQueueScratch(scratch); }); + MAKE_NAMED_SCOPE_GUARD(scratchGuard, [&]() { ReleaseQueueMainScratch(scratch); }); - if (scratch.size != 0) { - AcquireQueueScratch(scratch); - if (scratch.queue_base == nullptr) { + if (scratch.main_size != 0) { + AcquireQueueMainScratch(scratch); + if (scratch.main_queue_base == nullptr) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } } @@ -1598,18 +1598,19 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, return HSA_STATUS_SUCCESS; } -void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { - assert(scratch.queue_base == nullptr && "AcquireQueueScratch called while holding scratch."); +void GpuAgent::AcquireQueueMainScratch(ScratchInfo& scratch) { + assert(scratch.main_queue_base == nullptr && + "AcquireQueueMainScratch called while holding scratch."); bool need_queue_scratch_base = (isa_->GetMajorVersion() > 8); - if (scratch.size == 0) { - scratch.size = queue_scratch_len_; - scratch.size_per_thread = scratch_per_thread_; + if (scratch.main_size == 0) { + scratch.main_size = queue_scratch_len_; + scratch.main_size_per_thread = scratch_per_thread_; } scratch.retry = false; // Fail scratch allocation if per wave limits are exceeded. - uint64_t size_per_wave = AlignUp(scratch.size_per_thread * properties_.WaveFrontSize, 1024); + uint64_t size_per_wave = AlignUp(scratch.main_size_per_thread * properties_.WaveFrontSize, 1024); if (size_per_wave > MAX_WAVE_SCRATCH) return; /* @@ -1636,14 +1637,16 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { Limit total bound small scratch allocations to 1/8th of scratch pool and 1/4 of that for a single allocation. */ + bool large; + ScopedAcquire lock(&scratch_lock_); - size_t small_limit = scratch_pool_.size() >> 3; - const size_t single_scratch_limit = - core::Runtime::runtime_singleton_->flag().scratch_single_limit(); + const size_t small_limit = scratch_pool_.size() >> 3; bool use_reclaim = true; - bool large = (scratch.size > single_scratch_limit) || + + large = (scratch.main_size > scratch.use_once_limit) || ((scratch_pool_.size() - scratch_pool_.remaining() - scratch_cache_.free_bytes() + - scratch.size) > small_limit); + scratch.main_size) > small_limit); + if ((isa_->GetMajorVersion() < 8) || core::Runtime::runtime_singleton_->flag().no_scratch_reclaim()) { large = false; @@ -1652,10 +1655,10 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // If large is selected then the scratch will not be retained. // In that case allocate the minimum necessary for the dispatch since we don't need all slots. - if (large) scratch.size = scratch.dispatch_size; + if (large) scratch.main_size = scratch.dispatch_size; // Ensure mapping will be in whole pages. - scratch.size = AlignUp(scratch.size, 4096); + scratch.main_size = AlignUp(scratch.main_size, 4096); /* Sequence of attempts is: @@ -1673,31 +1676,32 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { [&]() { // Check scratch cache scratch.large = large; - if (scratch_cache_.alloc(scratch)) return; + if (scratch_cache_.allocMain(scratch)) return; // Attempt new allocation. for (int i = 0; i < 3; i++) { if (large) - scratch.queue_base = scratch_pool_.alloc_high(scratch.size); + scratch.main_queue_base = scratch_pool_.alloc_high(scratch.main_size); else - scratch.queue_base = scratch_pool_.alloc(scratch.size); - scratch.large = large | (scratch.queue_base > scratch_pool_.high_split()); + scratch.main_queue_base = scratch_pool_.alloc(scratch.main_size); + + scratch.large = large | (scratch.main_queue_base > scratch_pool_.high_split()); assert(((!scratch.large) | use_reclaim) && "Large scratch used with reclaim disabled."); - if (scratch.queue_base != nullptr) { + if (scratch.main_queue_base != nullptr) { HSAuint64 alternate_va; if ((profile_ == HSA_PROFILE_FULL) || - (hsaKmtMapMemoryToGPU(scratch.queue_base, scratch.size, &alternate_va) == + (hsaKmtMapMemoryToGPU(scratch.main_queue_base, scratch.main_size, &alternate_va) == HSAKMT_STATUS_SUCCESS)) { - if (scratch.large) scratch_used_large_ += scratch.size; - scratch_cache_.insert(scratch); + if (scratch.large) scratch_used_large_ += scratch.main_size; + scratch_cache_.insertMain(scratch); return; } } // Scratch request failed allocation or mapping. - scratch_pool_.free(scratch.queue_base); - scratch.queue_base = nullptr; + scratch_pool_.free(scratch.main_queue_base); + scratch.main_queue_base = nullptr; // Release cached scratch and retry. // First iteration trims unused blocks, second trims all. 3rd uses reserved memory @@ -1727,12 +1731,12 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // 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 (%ld) - reducing queue occupancy.\n", - scratch.size); + scratch.main_size); const uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU; const uint64_t se_per_xcc = properties_.NumShaderBanks / properties_.NumXcc; - const uint64_t total_waves = scratch.size / size_per_wave; - uint64_t waves_per_cu = AlignUp(total_waves / num_cus, scratch.waves_per_group); + const uint64_t total_waves = scratch.main_size / size_per_wave; + uint64_t waves_per_cu = AlignUp(total_waves / num_cus, scratch.main_waves_per_group); while (waves_per_cu != 0) { size_t size = waves_per_cu * num_cus * size_per_wave; @@ -1742,14 +1746,14 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { ((profile_ == HSA_PROFILE_FULL) || (hsaKmtMapMemoryToGPU(base, size, &alternate_va) == HSAKMT_STATUS_SUCCESS))) { // Scratch allocated and either full profile or map succeeded. - scratch.queue_base = base; - scratch.size = size; + scratch.main_queue_base = base; + scratch.main_size = size; scratch.large = true; - scratch_used_large_ += scratch.size; - scratch_cache_.insert(scratch); + scratch_used_large_ += scratch.main_size; + scratch_cache_.insertMain(scratch); 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); + debug_print(" %ld scratch mapped, %.2f%% occupancy.\n", scratch.main_size, + float(waves_per_cu * num_cus) / scratch.dispatch_slots * 100.0f); return; } scratch_pool_.free(base); @@ -1757,30 +1761,31 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { // Wave count must be divisible by #SEs in an XCC. If occupancy must be reduced // such that waves_per_cu < waves_per_group, continue reducing by #SEs per XCC // (only allowed if waves_per_group is a multiple #SEs per XCC). - waves_per_cu -= (waves_per_cu <= scratch.waves_per_group && - se_per_xcc < scratch.waves_per_group && - scratch.waves_per_group % se_per_xcc == 0) ? - se_per_xcc : scratch.waves_per_group; + waves_per_cu -= (waves_per_cu <= scratch.main_waves_per_group && + se_per_xcc < scratch.main_waves_per_group && + scratch.main_waves_per_group % se_per_xcc == 0) + ? se_per_xcc + : scratch.main_waves_per_group; } // Failed to allocate minimal scratch - assert(scratch.queue_base == nullptr && "bad scratch data"); + assert(scratch.main_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"); return; }(); - scratch.queue_process_offset = need_queue_scratch_base - ? uintptr_t(scratch.queue_base) - : uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base()); + scratch.main_queue_process_offset = need_queue_scratch_base + ? uintptr_t(scratch.main_queue_base) + : uintptr_t(scratch.main_queue_base) - uintptr_t(scratch_pool_.base()); } -void GpuAgent::ReleaseQueueScratch(ScratchInfo& scratch) { - if (scratch.queue_base == nullptr) return; - +void GpuAgent::ReleaseQueueMainScratch(ScratchInfo& scratch) { ScopedAcquire lock(&scratch_lock_); - scratch_cache_.free(scratch); - scratch.queue_base = nullptr; + if (scratch.main_queue_base == nullptr) return; + + scratch_cache_.freeMain(scratch); + scratch.main_queue_base = nullptr; } void GpuAgent::ReleaseScratch(void* base, size_t size, bool large) {