Correct scratch retry logic.

scratch_used_large_ was uninitialized leading to the observed hang.
DynamicScratchHandler would wait for a large scratch release despite no
large scratch having yet been allocated.  Fixes .

The patch also removes a potential race between AddScratchNotifier and
ReleaseQueueScratch.  The race condition does not exist today since both
scratch alloc and release run on the same thread.  The changes will
prevent this potential race from manifesting if the async event handler
is ever updated to use multiple threads.

Also enhances scratch occupancy reduction reporting.  Reporting now
prints the initial request size as well as the allocated size and the
effect on occupancy this has.  Occupancy is computed in terms of the
requesting dispatch grid size so may be >100%.

Change-Id: I0fc5ee01467ff4c29bdd25d545177c97862c3bd9

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

@@ -71,6 +71,8 @@ struct ScratchInfo {
   ptrdiff_t queue_process_offset;
   bool large;
   bool retry;
+  hsa_signal_t queue_retry;
+  uint64_t wanted_slots;
 };
 
 // @brief Interface to represent a GPU agent.
@@ -265,22 +267,8 @@ class GpuAgent : public GpuAgentInt {
   // @brief Override from amd::GpuAgentInt.
   void ReleaseQueueScratch(ScratchInfo& scratch) override;
 
-  // @brief Register signal for notification when scratch may become available.
-  // @p signal is notified by OR'ing with @p value.
-  void AddScratchNotifier(hsa_signal_t signal, hsa_signal_value_t value) {
-    ScopedAcquire<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,
-                     hsa_amd_profiling_dispatch_time_t& time) override;
+  void TranslateTime(core::Signal* signal, hsa_amd_profiling_dispatch_time_t& time) override;
 
   // @brief Override from amd::GpuAgentInt.
   void TranslateTime(core::Signal* signal, hsa_amd_profiling_async_copy_time_t& time) override;
@@ -494,6 +482,16 @@ class GpuAgent : public GpuAgentInt {
   // @brief Setup GWS accessing queue.
   void InitGWS();
 
+  // @brief Register signal for notification when scratch may become available.
+  // @p signal is notified by OR'ing with @p value.
+  bool AddScratchNotifier(hsa_signal_t signal, hsa_signal_value_t value) {
+    if (signal.handle != 0) return false;
+    scratch_notifiers_[signal] = value;
+    return true;
+  }
+
+  // @brief Deregister scratch notification signals.
+  void ClearScratchNotifiers() { scratch_notifiers_.clear(); }
 
   // Bind index of peer device that is connected via xGMI links
   lazy_ptr<core::Blit>& GetXgmiBlit(const core::Agent& peer_agent);

runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp

@@ -189,9 +189,6 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
                    0);
 #endif
 
-  // Initialize scratch memory related entities
-  InitScratchSRD();
-
   // Set group and private memory apertures in amd_queue_.
   auto& regions = agent->regions();
 
@@ -257,6 +254,11 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
     assert(Signal != nullptr && "Should have thrown!\n");
     amd_queue_.queue_inactive_signal = core::DefaultSignal::Convert(Signal);
   }
+
+  // Initialize scratch memory related entities
+  queue_scratch_.queue_retry = amd_queue_.queue_inactive_signal;
+  InitScratchSRD();
+
   if (AMD::hsa_amd_signal_async_handler(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE,
                                         0, DynamicScratchHandler, this) != HSA_STATUS_SUCCESS)
     throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
@@ -729,7 +731,6 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) {
 
   if ((queue->dynamicScratchState & ERROR_HANDLER_SCRATCH_RETRY) == ERROR_HANDLER_SCRATCH_RETRY) {
     queue->dynamicScratchState &= ~ERROR_HANDLER_SCRATCH_RETRY;
-    queue->agent_->RemoveScratchNotifier(queue->amd_queue_.queue_inactive_signal);
     changeWait = true;
     waitVal = 0;
     HSA::hsa_signal_and_relaxed(queue->amd_queue_.queue_inactive_signal, ~0x8000000000000000ull);
@@ -771,18 +772,23 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) {
 
       uint32_t scratch_request = pkt.dispatch.private_segment_size;
 
+      const uint32_t MaxScratchSlots =
+          (queue->amd_queue_.max_cu_id + 1) * queue->agent_->properties().MaxSlotsScratchCU;
+
       scratch.size_per_thread = scratch_request;
       scratch.lanes_per_wave = (error_code & 0x400) ? 32 : 64;
       // Align whole waves to 1KB.
       scratch.size_per_thread = AlignUp(scratch.size_per_thread, 1024 / scratch.lanes_per_wave);
-      scratch.size = scratch.size_per_thread * (queue->amd_queue_.max_cu_id + 1) *
-          queue->agent_->properties().MaxSlotsScratchCU * scratch.lanes_per_wave;
+      scratch.size = scratch.size_per_thread * MaxScratchSlots * scratch.lanes_per_wave;
+#ifndef NDEBUG
+      scratch.wanted_slots = ((uint64_t(pkt.dispatch.grid_size_x) * pkt.dispatch.grid_size_y) *
+                              pkt.dispatch.grid_size_z) / scratch.lanes_per_wave;
+      scratch.wanted_slots = Min(scratch.wanted_slots, MaxScratchSlots);
+#endif
 
       queue->agent_->AcquireQueueScratch(scratch);
 
       if (scratch.retry) {
-        queue->agent_->AddScratchNotifier(queue->amd_queue_.queue_inactive_signal,
-                                          0x8000000000000000ull);
         queue->dynamicScratchState |= ERROR_HANDLER_SCRATCH_RETRY;
         changeWait = true;
         waitVal = error_code;

runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp

@@ -77,6 +77,7 @@ GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props)
     : GpuAgentInt(node),
       properties_(node_props),
       current_coherency_type_(HSA_AMD_COHERENCY_TYPE_COHERENT),
+      scratch_used_large_(0),
       queues_(),
       local_region_(NULL),
       is_kv_device_(false),
@@ -936,7 +937,7 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
   }
 
   // Allocate scratch memory
-  ScratchInfo scratch;
+  ScratchInfo scratch = {0};
   if (private_segment_size == UINT_MAX) {
     private_segment_size = (profile_ == HSA_PROFILE_BASE) ? 0 : scratch_per_thread_;
   }
@@ -1038,15 +1039,16 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) {
 
   // Retry if large may yield needed space.
   if (scratch_used_large_ != 0) {
-    scratch.retry = true;
+    if (AddScratchNotifier(scratch.queue_retry, 0x8000000000000000ull)) scratch.retry = true;
     return;
   }
 
   // Attempt to trim the maximum number of concurrent waves to allow scratch to fit.
   if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message())
-    debug_print("Failed to map requested scratch - reducing queue occupancy.\n");
-  uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU;
-  uint64_t total_waves = scratch.size / size_per_wave;
+    debug_print("Failed to map requested scratch (%ld) - reducing queue occupancy.\n",
+                scratch.size);
+  const uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU;
+  const uint64_t total_waves = scratch.size / size_per_wave;
   uint64_t waves_per_cu = total_waves / num_cus;
   while (waves_per_cu != 0) {
     size_t size = waves_per_cu * num_cus * size_per_wave;
@@ -1058,12 +1060,14 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) {
       // Scratch allocated and either full profile or map succeeded.
       scratch.queue_base = base;
       scratch.size = size;
-      scratch.queue_process_offset =
-          (need_queue_scratch_base)
-              ? uintptr_t(scratch.queue_base)
-              : uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base());
+      scratch.queue_process_offset = (need_queue_scratch_base)
+          ? uintptr_t(scratch.queue_base)
+          : uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base());
       scratch.large = true;
       scratch_used_large_ += scratch.size;
+      if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message())
+        debug_print("  %ld scratch mapped, %.2f%% occupancy.\n", scratch.size,
+                    float(waves_per_cu * num_cus) / scratch.wanted_slots * 100.0f);
       return;
     }
     scratch_pool_.free(base);
@@ -1073,7 +1077,7 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) {
   // Failed to allocate minimal scratch
   assert(scratch.queue_base == nullptr && "bad scratch data");
   if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message())
-    debug_print("Could not allocate scratch for one wave per CU.\n");
+    debug_print("  Could not allocate scratch for one wave per CU.\n");
 }
 
 void GpuAgent::ReleaseQueueScratch(ScratchInfo& scratch) {
@@ -1093,12 +1097,13 @@ void GpuAgent::ReleaseQueueScratch(ScratchInfo& scratch) {
   if (scratch.large) scratch_used_large_ -= scratch.size;
 
   // Notify waiters that additional scratch may be available.
-  for (auto notifier : scratch_notifiers_)
+  for (auto notifier : scratch_notifiers_) {
     HSA::hsa_signal_or_relaxed(notifier.first, notifier.second);
+  }
+  ClearScratchNotifiers();
 }
 
-void GpuAgent::TranslateTime(core::Signal* signal,
-                             hsa_amd_profiling_dispatch_time_t& time) {
+void GpuAgent::TranslateTime(core::Signal* signal, hsa_amd_profiling_dispatch_time_t& time) {
   uint64_t start, end;
   signal->GetRawTs(false, start, end);
   // Order is important, we want to translate the end time first to ensure that packet duration is