From 7caf9633f6e3cd7bb57548a7d01cb2b9961418e1 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Wed, 28 Mar 2018 11:11:08 -0500 Subject: [PATCH] Support large scratch allocations and reclaim. Also improve small_heap used for scratch region allocation. Change-Id: Ib7311b663b38968d88ebc355b81e12c0863dc541 --- runtime/hsa-runtime/core/inc/amd_aql_queue.h | 2 +- runtime/hsa-runtime/core/inc/amd_gpu_agent.h | 32 ++- runtime/hsa-runtime/core/inc/signal.h | 13 ++ .../core/runtime/amd_aql_queue.cpp | 78 ++++++-- .../core/runtime/amd_gpu_agent.cpp | 58 ++++-- runtime/hsa-runtime/core/util/small_heap.cpp | 185 +++++++++--------- runtime/hsa-runtime/core/util/small_heap.h | 83 ++++---- runtime/hsa-runtime/inc/amd_hsa_common.h | 8 +- runtime/hsa-runtime/inc/amd_hsa_queue.h | 3 +- 9 files changed, 296 insertions(+), 166 deletions(-) diff --git a/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/runtime/hsa-runtime/core/inc/amd_aql_queue.h index a46471c820..795928bfb1 100644 --- a/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -252,7 +252,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo // Error handler control variable. std::atomic dynamicScratchState; - enum { ERROR_HANDLER_DONE = 1, ERROR_HANDLER_TERMINATE = 2 }; + enum { ERROR_HANDLER_DONE = 1, ERROR_HANDLER_TERMINATE = 2, ERROR_HANDLER_SCRATCH_RETRY = 4 }; // Shared event used for queue errors static HsaEvent* queue_event_; diff --git a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index fd5750eace..2de2fd9a61 100644 --- a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -46,6 +46,7 @@ #define HSA_RUNTIME_CORE_INC_AMD_GPU_AGENT_H_ #include +#include #include "hsakmt.h" @@ -67,6 +68,8 @@ struct ScratchInfo { size_t size; size_t size_per_thread; ptrdiff_t queue_process_offset; + bool large; + bool retry; }; // @brief Interface to represent a GPU agent. @@ -103,15 +106,15 @@ class GpuAgentInt : public core::Agent { // @brief Carve scratch memory from scratch pool. // - // @param [out] scratch Structure to be populated with the carved memory + // @param [in/out] scratch Structure to be populated with the carved memory // information. virtual void AcquireQueueScratch(ScratchInfo& scratch) = 0; // @brief Release scratch memory back to scratch pool. // - // @param [in] base Address of scratch memory previously acquired with - // call to ::AcquireQueueScratch. - virtual void ReleaseQueueScratch(void* base) = 0; + // @param [in/out] scratch Scratch memory previously acquired with call to + // ::AcquireQueueScratch. + virtual void ReleaseQueueScratch(ScratchInfo& base) = 0; // @brief Translate the kernel start and end dispatch timestamp from agent // domain to host domain. @@ -257,7 +260,20 @@ class GpuAgent : public GpuAgentInt { void AcquireQueueScratch(ScratchInfo& scratch) override; // @brief Override from amd::GpuAgentInt. - void ReleaseQueueScratch(void* base) override; + 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, @@ -368,6 +384,12 @@ class GpuAgent : public GpuAgentInt { // @brief Object to manage scratch memory. SmallHeap scratch_pool_; + // @brief Current short duration scratch memory size. + size_t scratch_used_large_; + + // @brief Notifications for scratch release. + std::map scratch_notifiers_; + // @brief Default scratch size per queue. size_t queue_scratch_len_; diff --git a/runtime/hsa-runtime/core/inc/signal.h b/runtime/hsa-runtime/core/inc/signal.h index 9c0d25618c..7f676d4790 100644 --- a/runtime/hsa-runtime/core/inc/signal.h +++ b/runtime/hsa-runtime/core/inc/signal.h @@ -46,6 +46,7 @@ #define HSA_RUNTME_CORE_INC_SIGNAL_H_ #include +#include #include "hsakmt.h" @@ -60,6 +61,18 @@ #include "inc/amd_hsa_signal.h" +// Allow hsa_signal_t to be keys in STL structures. +namespace std { +template <> struct less { + __forceinline bool operator()(const hsa_signal_t& x, const hsa_signal_t& y) const { + return x.handle < y.handle; + } + typedef hsa_signal_t first_argument_type; + typedef hsa_signal_t second_argument_type; + typedef bool result_type; +}; +} + namespace core { class Agent; class Signal; diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 9b72bfc3f9..f96c433163 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -291,8 +291,8 @@ AqlQueue::~AqlQueue() { } Inactivate(); + agent_->ReleaseQueueScratch(queue_scratch_); FreeRegisteredRingBuffer(); - agent_->ReleaseQueueScratch(queue_scratch_.queue_base); HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal); if (core::g_use_interrupt_wait) { ScopedAcquire lock(&queue_lock_); @@ -704,24 +704,54 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { AqlQueue* queue = (AqlQueue*)arg; hsa_status_t errorCode = HSA_STATUS_SUCCESS; bool fatal = false; + bool changeWait = false; + hsa_signal_value_t waitVal; + + 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); + error_code &= ~0x8000000000000000ull; + } // Process errors only if queue is not terminating. if ((queue->dynamicScratchState & ERROR_HANDLER_TERMINATE) != ERROR_HANDLER_TERMINATE) { - // Process only one queue error, don't fall through. + 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->InitScratchSRD(); + + HSA::hsa_signal_store_relaxed(queue->amd_queue_.queue_inactive_signal, 0); + // Resumes queue processing. + atomic::Store(&queue->amd_queue_.queue_properties, + queue->amd_queue_.queue_properties & (~AMD_QUEUE_PROPERTIES_USE_SCRATCH_ONCE), + std::memory_order_release); + atomic::Fence(std::memory_order_release); + return true; + } + + // Process only one queue error. if (error_code == 1) { // Insufficient scratch - recoverable, don't process dynamic scratch if errors are present. auto& scratch = queue->queue_scratch_; - queue->agent_->ReleaseQueueScratch(scratch.queue_base); + queue->agent_->ReleaseQueueScratch(scratch); - uint64_t pkt_slot_idx = queue->amd_queue_.read_dispatch_id % queue->amd_queue_.hsa_queue.size; + uint64_t pkt_slot_idx = + queue->amd_queue_.read_dispatch_id & (queue->amd_queue_.hsa_queue.size - 1); const core::AqlPacket& pkt = ((core::AqlPacket*)queue->amd_queue_.hsa_queue.base_address)[pkt_slot_idx]; uint32_t scratch_request = pkt.dispatch.private_segment_size; - scratch.size_per_thread = Max(uint32_t(scratch.size_per_thread * 2), scratch_request); + scratch.size_per_thread = scratch_request; // Align whole waves to 1KB. scratch.size_per_thread = AlignUp(scratch.size_per_thread, 16); scratch.size = scratch.size_per_thread * (queue->amd_queue_.max_cu_id + 1) * @@ -729,11 +759,26 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { queue->agent_->AcquireQueueScratch(scratch); - // Out of scratch - promote error - if (scratch.queue_base == NULL) errorCode = HSA_STATUS_ERROR_OUT_OF_RESOURCES; - - // Reset scratch memory related entities for the queue - queue->InitScratchSRD(); + if (scratch.retry) { + queue->agent_->AddScratchNotifier(queue->amd_queue_.queue_inactive_signal, + 0x8000000000000000ull); + queue->dynamicScratchState |= ERROR_HANDLER_SCRATCH_RETRY; + changeWait = true; + waitVal = error_code; + } else { + // Out of scratch - promote error + if (scratch.queue_base == nullptr) { + errorCode = HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } else { + // Mark large scratch allocation for single use. + if (scratch.large) + queue->amd_queue_.queue_properties |= AMD_QUEUE_PROPERTIES_USE_SCRATCH_ONCE; + // Reset scratch memory related entities for the queue + queue->InitScratchSRD(); + // Restart the queue. + HSA::hsa_signal_store_screlease(queue->amd_queue_.queue_inactive_signal, 0); + } + } } else if ((error_code & 2) == 2) { // Invalid dim errorCode = HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS; @@ -765,7 +810,12 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { } if (errorCode == HSA_STATUS_SUCCESS) { - HSA::hsa_signal_store_relaxed(queue->amd_queue_.queue_inactive_signal, 0); + if (changeWait) { + core::Runtime::runtime_singleton_->SetAsyncSignalHandler( + queue->amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE, waitVal, + DynamicScratchHandler, queue); + return false; + } return true; } @@ -774,9 +824,9 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { queue->errors_callback_(errorCode, queue->public_handle(), queue->errors_data_); } if (fatal) { - //Temporarilly removed until there is clarity on exactly what debugtrap's semantics are. - //assert(false && "Fatal queue error"); - //std::abort(); + // Temporarilly removed until there is clarity on exactly what debugtrap's semantics are. + // assert(false && "Fatal queue error"); + // std::abort(); } } // Copy here is to protect against queue being released between setting the scratch state and diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index 83a69520e3..6421eb06e7 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -330,7 +330,6 @@ void GpuAgent::InitScratchPool() { // scratch/thread const uint32_t num_cu = properties_.NumFComputeCores / properties_.NumSIMDPerCU; - queue_scratch_len_ = 0; queue_scratch_len_ = AlignUp(32 * 64 * num_cu * scratch_per_thread_, 65536); size_t max_scratch_len = queue_scratch_len_ * max_queues_; @@ -352,7 +351,7 @@ void GpuAgent::InitScratchPool() { if (HSAKMT_STATUS_SUCCESS == err) { new (&scratch_pool_) SmallHeap(scratch_base, max_scratch_len); } else { - new (&scratch_pool_) SmallHeap(NULL, 0); + new (&scratch_pool_) SmallHeap(); } } @@ -892,8 +891,9 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, const uint32_t num_cu = properties_.NumFComputeCores / properties_.NumSIMDPerCU; scratch.size = scratch.size_per_thread * 32 * 64 * num_cu; scratch.queue_base = nullptr; + scratch.queue_process_offset = 0; - MAKE_NAMED_SCOPE_GUARD(scratchGuard, [&]() { ReleaseQueueScratch(scratch.queue_base); }); + MAKE_NAMED_SCOPE_GUARD(scratchGuard, [&]() { ReleaseQueueScratch(scratch); }); if (scratch.size != 0) { AcquireQueueScratch(scratch); @@ -921,30 +921,46 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { scratch.size_per_thread = scratch_per_thread_; } + scratch.retry = false; + ScopedAcquire lock(&scratch_lock_); - scratch.queue_base = scratch_pool_.alloc(scratch.size); + bool large = (scratch.size > 6 * 1024 * 1024) || + (scratch_pool_.size() - scratch_pool_.remaining() > 24 * 6 * 1024 * 1024); + if (large) + scratch.queue_base = scratch_pool_.alloc_high(scratch.size); + else + scratch.queue_base = scratch_pool_.alloc(scratch.size); + large |= scratch.queue_base > scratch_pool_.high_split(); + scratch.large = large; + scratch.queue_process_offset = (need_queue_scratch_base) ? uintptr_t(scratch.queue_base) : uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base()); - if (scratch.queue_base != NULL) { + if (scratch.queue_base != nullptr) { if (profile_ == HSA_PROFILE_FULL) return; if (profile_ == HSA_PROFILE_BASE) { HSAuint64 alternate_va; - if (HSAKMT_STATUS_SUCCESS == - hsaKmtMapMemoryToGPU(scratch.queue_base, scratch.size, &alternate_va)) + if (hsaKmtMapMemoryToGPU(scratch.queue_base, scratch.size, &alternate_va) == + HSAKMT_STATUS_SUCCESS) { + if (large) scratch_used_large_ += scratch.size; return; + } } } // Scratch request failed allocation or mapping. scratch_pool_.free(scratch.queue_base); - scratch.queue_base = NULL; + scratch.queue_base = nullptr; -// 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. + // Retry if large may yield needed space. + if (scratch_used_large_ != 0) { + 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; @@ -955,7 +971,7 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { size_t size = waves_per_cu * num_cus * size_per_wave; void* base = scratch_pool_.alloc(size); HSAuint64 alternate_va; - if ((base != NULL) && + if ((base != nullptr) && ((profile_ == HSA_PROFILE_FULL) || (hsaKmtMapMemoryToGPU(base, size, &alternate_va) == HSAKMT_STATUS_SUCCESS))) { // Scratch allocated and either full profile or map succeeded. @@ -965,6 +981,8 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { (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; return; } scratch_pool_.free(base); @@ -972,23 +990,29 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { } // Failed to allocate minimal scratch - assert(scratch.queue_base == NULL && "bad scratch data"); + 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"); } -void GpuAgent::ReleaseQueueScratch(void* base) { - if (base == NULL) { +void GpuAgent::ReleaseQueueScratch(ScratchInfo& scratch) { + if (scratch.queue_base == nullptr) { return; } ScopedAcquire lock(&scratch_lock_); if (profile_ == HSA_PROFILE_BASE) { - if (HSAKMT_STATUS_SUCCESS != hsaKmtUnmapMemoryToGPU(base)) { + if (HSAKMT_STATUS_SUCCESS != hsaKmtUnmapMemoryToGPU(scratch.queue_base)) { assert(false && "Unmap scratch subrange failed!"); } } - scratch_pool_.free(base); + scratch_pool_.free(scratch.queue_base); + + if (scratch.large) scratch_used_large_ -= scratch.size; + + // Notify waiters that additional scratch may be available. + for (auto notifier : scratch_notifiers_) + HSA::hsa_signal_or_relaxed(notifier.first, notifier.second); } void GpuAgent::TranslateTime(core::Signal* signal, diff --git a/runtime/hsa-runtime/core/util/small_heap.cpp b/runtime/hsa-runtime/core/util/small_heap.cpp index 9b8998bf30..6cd8e117dc 100644 --- a/runtime/hsa-runtime/core/util/small_heap.cpp +++ b/runtime/hsa-runtime/core/util/small_heap.cpp @@ -42,25 +42,47 @@ #include "small_heap.h" -SmallHeap::memory_t::iterator SmallHeap::merge( - SmallHeap::memory_t::iterator& keep, - SmallHeap::memory_t::iterator& destroy) { - assert((char*)keep->first + keep->second.len == (char*)destroy->first && - "Invalid merge"); - assert(keep->second.isfree() && "Merge with allocated block"); - assert(destroy->second.isfree() && "Merge with allocated block"); +// Inserts node into freelist after place. +// Assumes node will not be an end of the list (list has guard nodes). +void SmallHeap::insertafter(SmallHeap::iterator_t place, SmallHeap::iterator_t node) { + assert(place->first < node->first && "Order violation"); + assert(isfree(place->second) && "Freelist operation error."); + iterator_t next = place->second.next; + node->second.next = next; + node->second.prior = place; + place->second.next = node; + next->second.prior = node; +} - keep->second.len += destroy->second.len; - keep->second.next_free = destroy->second.next_free; - if (!destroy->second.islastfree()) - memory[destroy->second.next_free].prior_free = keep->first; +// Removes node from freelist. +// Assumes node will not be an end of the list (list has guard nodes). +void SmallHeap::remove(SmallHeap::iterator_t node) { + assert(isfree(node->second) && "Freelist operation error."); + node->second.prior->second.next = node->second.next; + node->second.next->second.prior = node->second.prior; + setused(node->second); +} - memory.erase(destroy); - return keep; +// Returns high if merge failed or the merged node. +SmallHeap::memory_t::iterator SmallHeap::merge(SmallHeap::memory_t::iterator low, + SmallHeap::memory_t::iterator high) { + assert(isfree(low->second) && "Merge with allocated block"); + assert(isfree(high->second) && "Merge with allocated block"); + + if ((char*)low->first + low->second.len != (char*)high->first) return high; + + assert(!islastfree(high->second) && "Illegal merge."); + + low->second.len += high->second.len; + low->second.next = high->second.next; + high->second.next->second.prior = low; + + memory.erase(high); + return low; } void SmallHeap::free(void* ptr) { - if (ptr == NULL) return; + if (ptr == nullptr) return; auto iterator = memory.find(ptr); @@ -70,105 +92,90 @@ void SmallHeap::free(void* ptr) { return; } - const auto start_guard = memory.find(0); - const auto end_guard = memory.find((void*)0xFFFFFFFFFFFFFFFFull); - // Return memory to total and link node into free list total_free += iterator->second.len; - if (first_free < iterator->first) { - auto before = iterator; - before--; - while (before != start_guard && !before->second.isfree()) before--; - assert(before->second.next_free > iterator->first && - "Inconsistency in small heap."); - iterator->second.prior_free = before->first; - iterator->second.next_free = before->second.next_free; - before->second.next_free = iterator->first; - if (!iterator->second.islastfree()) - memory[iterator->second.next_free].prior_free = iterator->first; - } else { - iterator->second.setfirstfree(); - iterator->second.next_free = first_free; - first_free = iterator->first; - if (!iterator->second.islastfree()) - memory[iterator->second.next_free].prior_free = iterator->first; - } - // Attempt compaction + // Could also traverse the free list which might be faster in some cases. auto before = iterator; before--; - if (before != start_guard) { - if (before->second.isfree()) { - iterator = merge(before, iterator); - } - } + while (!isfree(before->second)) before--; + assert(before->second.next->first > iterator->first && "Inconsistency in small heap."); + insertafter(before, iterator); - auto after = iterator; - after++; - if (after != end_guard) { - if (after->second.isfree()) { - iterator = merge(iterator, after); - } - } + // Attempt compaction + iterator = merge(before, iterator); + merge(iterator, iterator->second.next); + + // Update lowHighBondary + high.erase(ptr); } void* SmallHeap::alloc(size_t bytes) { // Is enough memory available? - if ((bytes > total_free) || (bytes == 0)) return NULL; + if ((bytes > total_free) || (bytes == 0)) return nullptr; - memory_t::iterator current; - memory_t::iterator prior; + iterator_t current; // Walk the free list and allocate at first fitting location - prior = current = memory.find(first_free); - while (true) { + current = firstfree(); + while (!islastfree(current->second)) { if (bytes <= current->second.len) { // Decrement from total total_free -= bytes; - // Is allocation an exact fit? - if (bytes == current->second.len) { - if (prior == current) { - first_free = current->second.next_free; - if (!current->second.islastfree()) - memory[current->second.next_free].setfirstfree(); - } else { - prior->second.next_free = current->second.next_free; - if (!current->second.islastfree()) - memory[current->second.next_free].prior_free = prior->first; - } - current->second.next_free = NULL; - return current->first; - } else { - // Split current node + // Split node + if (bytes != current->second.len) { void* remaining = (char*)current->first + bytes; Node& node = memory[remaining]; - node.next_free = current->second.next_free; - node.prior_free = current->second.prior_free; node.len = current->second.len - bytes; current->second.len = bytes; - - if (prior == current) { - first_free = remaining; - node.setfirstfree(); - } else { - prior->second.next_free = remaining; - node.prior_free = prior->first; - } - if (!node.islastfree()) memory[node.next_free].prior_free = remaining; - - current->second.next_free = NULL; - return current->first; + insertafter(current, memory.find(remaining)); } + + remove(current); + return current->first; } - - // End of free list? - if (current->second.islastfree()) break; - - prior = current; - current = memory.find(current->second.next_free); + current = current->second.next; } + assert(current->second.len == 0 && "Freelist corruption."); // Can't service the request due to fragmentation - return NULL; + return nullptr; +} + +void* SmallHeap::alloc_high(size_t bytes) { + // Is enough memory available? + if ((bytes > total_free) || (bytes == 0)) return nullptr; + + iterator_t current; + + // Walk the free list and allocate at first fitting location + current = lastfree(); + while (!isfirstfree(current->second)) { + if (bytes <= current->second.len) { + // Decrement from total + total_free -= bytes; + + void* alloc; + // Split node + if (bytes != current->second.len) { + alloc = (char*)current->first + current->second.len - bytes; + current->second.len -= bytes; + Node& node = memory[alloc]; + node.len = bytes; + setused(node); + } else { + alloc = current->first; + remove(current); + } + + high.insert(alloc); + return alloc; + } + current = current->second.prior; + } + assert(current->second.len == 0 && "Freelist corruption."); + + // Can't service the request due to fragmentation + return nullptr; } diff --git a/runtime/hsa-runtime/core/util/small_heap.h b/runtime/hsa-runtime/core/util/small_heap.h index fecfe857f5..d9064bba66 100644 --- a/runtime/hsa-runtime/core/util/small_heap.h +++ b/runtime/hsa-runtime/core/util/small_heap.h @@ -47,68 +47,81 @@ #ifndef HSA_RUNTME_CORE_UTIL_SMALL_HEAP_H_ #define HSA_RUNTME_CORE_UTIL_SMALL_HEAP_H_ -#include "utils.h" - #include +#include +#include "utils.h" + class SmallHeap { - public: - class Node { - public: - size_t len; - void* next_free; - void* prior_free; - static const intptr_t END = -1; + private: + struct Node; + typedef std::map memory_t; + typedef memory_t::iterator iterator_t; - __forceinline bool isfree() const { return next_free != NULL; } - __forceinline bool islastfree() const { return intptr_t(next_free) == END; } - __forceinline bool isfirstfree() const { - return intptr_t(prior_free) == END; - } - __forceinline void setlastfree() { - *reinterpret_cast(&next_free) = END; - } - __forceinline void setfirstfree() { - *reinterpret_cast(&prior_free) = END; - } + struct Node { + size_t len; + iterator_t next; + iterator_t prior; }; - private: - SmallHeap(const SmallHeap& rhs); - SmallHeap& operator=(const SmallHeap& rhs); + SmallHeap(const SmallHeap& rhs) = delete; + SmallHeap& operator=(const SmallHeap& rhs) = delete; void* const pool; const size_t length; size_t total_free; - void* first_free; - std::map memory; + memory_t memory; + std::set high; - typedef decltype(memory) memory_t; - memory_t::iterator merge(memory_t::iterator& keep, - memory_t::iterator& destroy); + __forceinline bool isfree(const Node& node) const { return node.next != memory.begin(); } + __forceinline bool islastfree(const Node& node) const { return node.next == memory.end(); } + __forceinline bool isfirstfree(const Node& node) const { return node.prior == memory.end(); } + __forceinline void setlastfree(Node& node) { node.next = memory.end(); } + __forceinline void setfirstfree(Node& node) { node.prior = memory.end(); } + __forceinline void setused(Node& node) { node.next = memory.begin(); } + + __forceinline iterator_t firstfree() { return memory.begin()->second.next; } + __forceinline iterator_t lastfree() { return memory.rbegin()->second.prior; } + void insertafter(iterator_t place, iterator_t node); + void remove(iterator_t node); + iterator_t merge(iterator_t low, iterator_t high); public: - SmallHeap() : pool(NULL), length(0), total_free(0) {} + SmallHeap() : pool(nullptr), length(0), total_free(0) {} SmallHeap(void* base, size_t length) : pool(base), length(length), total_free(length) { - first_free = pool; + assert(pool != nullptr && "Invalid base address."); + assert(pool != (void*)0xFFFFFFFFFFFFFFFFull && "Invalid base address."); + assert((char*)pool + length != (char*)0xFFFFFFFFFFFFFFFFull && "Invalid pool bounds."); + + Node& start = memory[0]; + Node& node = memory[pool]; + Node& end = memory[(void*)0xFFFFFFFFFFFFFFFFull]; + + start.len = 0; + start.next = memory.find(pool); + setfirstfree(start); - Node& node = memory[first_free]; node.len = length; - node.setlastfree(); - node.setfirstfree(); + node.prior = memory.begin(); + node.next = --memory.end(); - memory[0].len = 0; - memory[(void*)0xFFFFFFFFFFFFFFFFull].len = 0; + end.len = 0; + end.prior = start.next; + setlastfree(end); + + high.insert((void*)0xFFFFFFFFFFFFFFFFull); } void* alloc(size_t bytes); + void* alloc_high(size_t bytes); void free(void* ptr); void* base() const { return pool; } size_t size() const { return length; } size_t remaining() const { return total_free; } + void* high_split() const { return *high.begin(); } }; #endif diff --git a/runtime/hsa-runtime/inc/amd_hsa_common.h b/runtime/hsa-runtime/inc/amd_hsa_common.h index ca6a2b9838..bfb613ec47 100644 --- a/runtime/hsa-runtime/inc/amd_hsa_common.h +++ b/runtime/hsa-runtime/inc/amd_hsa_common.h @@ -75,8 +75,8 @@ // Creates enumeration entries for packed types. Enumeration entries include // bit shift amount, bit width, and bit mask. #define AMD_HSA_BITS_CREATE_ENUM_ENTRIES(name, shift, width) \ - name ## _SHIFT = (shift), \ - name ## _WIDTH = (width), \ + name##_SHIFT = (shift), \ + name##_WIDTH = (width), \ name = (((1 << (width)) - 1) << (shift)) \ // Gets bits for specified mask from specified src packed instance. @@ -85,7 +85,7 @@ // Sets val bits for specified mask in specified dst packed instance. #define AMD_HSA_BITS_SET(dst, mask, val) \ - dst &= (~(1 << mask ## _SHIFT) & ~mask); \ - dst |= (((val) << mask ## _SHIFT) & mask) \ + dst &= (~(1 << mask##_SHIFT) & ~mask); \ + dst |= (((val) << mask##_SHIFT) & mask) \ #endif // AMD_HSA_COMMON_H diff --git a/runtime/hsa-runtime/inc/amd_hsa_queue.h b/runtime/hsa-runtime/inc/amd_hsa_queue.h index b37bb53f36..2176e84706 100644 --- a/runtime/hsa-runtime/inc/amd_hsa_queue.h +++ b/runtime/hsa-runtime/inc/amd_hsa_queue.h @@ -53,7 +53,8 @@ enum amd_queue_properties_t { AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_QUEUE_PROPERTIES_IS_PTR64, 1, 1), AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_QUEUE_PROPERTIES_ENABLE_TRAP_HANDLER_DEBUG_SGPRS, 2, 1), AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, 3, 1), - AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_QUEUE_PROPERTIES_RESERVED1, 4, 28) + AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_QUEUE_PROPERTIES_USE_SCRATCH_ONCE, 4, 1), + AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_QUEUE_PROPERTIES_RESERVED1, 5, 27) }; // AMD Queue.