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

@@ -252,7 +252,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
 
   // Error handler control variable.
   std::atomic<uint32_t> 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_;

runtime/hsa-runtime/core/inc/amd_gpu_agent.h

@@ -46,6 +46,7 @@
 #define HSA_RUNTIME_CORE_INC_AMD_GPU_AGENT_H_
 
 #include <vector>
+#include <map>
 
 #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<KernelMutex> lock(&scratch_lock_);
+    scratch_notifiers_[signal] = value;
+  }
+
+  // @brief Deregister scratch notification signal.
+  void RemoveScratchNotifier(hsa_signal_t signal) {
+    ScopedAcquire<KernelMutex> 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<hsa_signal_t, hsa_signal_value_t> scratch_notifiers_;
+
   // @brief Default scratch size per queue.
   size_t queue_scratch_len_;
 

runtime/hsa-runtime/core/inc/signal.h

@@ -46,6 +46,7 @@
 #define HSA_RUNTME_CORE_INC_SIGNAL_H_
 
 #include <map>
+#include <functional>
 
 #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<hsa_signal_t> {
+  __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;

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<KernelMutex> 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

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<KernelMutex> 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<KernelMutex> 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,

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;
 }

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 <map>
+#include <set>
 
+#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<void*, Node> 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<intptr_t*>(&next_free) = END;
-    }
-    __forceinline void setfirstfree() {
-      *reinterpret_cast<intptr_t*>(&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<void*, Node> memory;
+  memory_t memory;
+  std::set<void*> 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

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

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.