From e68d671a51c79ed09bad83bc8d9fdfa0bf3c20b8 Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Tue, 25 Aug 2020 16:32:21 -0400
Subject: [PATCH] Reduce the default size of the signal pool

Implement dynamic signal pool grow per allocated queue

Change-Id: Ie8b17937d72c29cc49e59639c4b2023ea984b14c
---
 rocclr/device/rocm/rocvirtual.cpp | 22 ++++++++++++++++++++--
 rocclr/device/rocm/rocvirtual.hpp |  2 ++
 2 files changed, 22 insertions(+), 2 deletions(-)

diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp
index 2315655934..e1bf5a82ef 100644
--- a/rocclr/device/rocm/rocvirtual.cpp
+++ b/rocclr/device/rocm/rocvirtual.cpp
@@ -453,8 +453,17 @@ bool VirtualGPU::dispatchGenericAqlPacket(
 
   // TODO: placeholder to setup the kernel to populate start and end timestamp.
   if (timestamp_ != nullptr) {
+    if (current_signal_ >= signal_pool_.size()) {
+      ProfilingSignal profilingSignal = {};
+      if (HSA_STATUS_SUCCESS != hsa_signal_create(0, 0, nullptr, &profilingSignal.signal_)) {
+        LogPrintfError("Failed signal allocation id = %d", current_signal_);
+        return false;
+      }
+      signal_pool_.push_back(profilingSignal);
+      assert(current_signal_ < signal_pool_.size() && "Not enough signals");
+    }
     // Find signal slot
-    ProfilingSignal* profilingSignal = &signal_pool_[index & queueMask];
+    ProfilingSignal* profilingSignal = &signal_pool_[current_signal_++];
     // Make sure we save the old results in the TS structure
     if (profilingSignal->ts_ != nullptr) {
       profilingSignal->ts_->checkGpuTime();
@@ -526,6 +535,8 @@ bool VirtualGPU::dispatchGenericAqlPacket(
       LogPrintfError("Failed signal [0x%lx] wait", signal.handle);
       return false;
     }
+    // Reset the pool of signals
+    current_signal_ = 0;
   }
 
   return true;
@@ -634,6 +645,9 @@ bool VirtualGPU::releaseGpuMemoryFence() {
   // Release the pool, since runtime just completed a barrier
   resetKernArgPool();
 
+  // Reset the pool of signals
+  current_signal_ = 0;
+
   return true;
 }
 
@@ -757,8 +771,9 @@ bool VirtualGPU::create() {
   gpu_queue_ = roc_device_.acquireQueue(queue_size, cooperative_, cuMask_, priority_);
   if (!gpu_queue_) return false;
 
+  constexpr uint32_t kDefaultSignalPoolSize = 32;
   if (!initPool(dev().settings().kernargPoolSize_,
-                (profiling_ || (amd::IS_HIP)) ? queue_size : 0)) {
+                (profiling_ || (amd::IS_HIP)) ? kDefaultSignalPoolSize : 0)) {
     LogError("Couldn't allocate arguments/signals for the queue");
     return false;
   }
@@ -861,6 +876,9 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) {
       }
 
       resetKernArgPool();
+
+      // Reset the pool of signals
+      current_signal_ = 0;
     }
   } while (true);
 
diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp
index 862a2dd593..6ab868a659 100644
--- a/rocclr/device/rocm/rocvirtual.hpp
+++ b/rocclr/device/rocm/rocvirtual.hpp
@@ -356,6 +356,8 @@ class VirtualGPU : public device::VirtualDevice {
   uint kernarg_pool_cur_offset_;
 
   std::vector<ProfilingSignal> signal_pool_;  //!< Pool of signals for profiling
+  uint32_t current_signal_ = 0;               //!< Current avaialble signal in the pool
+
   friend class Timestamp;
 
   //  PM4 packet for gfx8 performance counter