AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Re-arrange and rename scratch elements that are used with main scratch

فهرست منبع 
Change-Id: I4c1ff8cf4121a06b586fe49c70400226506bf95e
This commit is contained in:
David Yat Sin
2023-11-15 16:34:19 +00:00
والد 0344c8c0b6
کامیت fa317f8c41
5فایلهای تغییر یافته به همراه163 افزوده شده و 139 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
runtime/hsa-runtime/core/inc/amd_gpu_agent.h
+7 -9
مشاهده پرونده
@@ -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;
runtime/hsa-runtime/core/inc/queue.h
+14
مشاهده پرونده
@@ -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);
runtime/hsa-runtime/core/inc/scratch_cache.h
+40 -36
مشاهده پرونده
@@ -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;
}
runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+40 -37
مشاهده پرونده
@@ -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;
}
runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+62 -57
مشاهده پرونده
@@ -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<KernelMutex> 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<KernelMutex> 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) {