SWDEV-268381 - Enable wait on CPU before SDMA transfer

- There is a performance regression with a HW wait for HSA signal
on ROCr async operation. For now move the logic back to CPU wait.

- Fix profiling issue with multiple HSA signal per single timestamp
object. Some copies require multiple ROCR calls and if profiling is
required, then the execution time is derived from all used signals.

Change-Id: Id003e4abb8c2de378eedc152a7e389500fc6f4ce

rocclr/device/rocm/rocblit.cpp

@@ -443,10 +443,12 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
     if (isSubwindowRectCopy ) {
       hsa_signal_t wait = gpu().Barriers().WaitSignal();
       hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp());
+      uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1;
+      hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait;
 
       // Copy memory line by line
       hsa_status_t status = hsa_amd_memory_async_copy_rect(&dstMem, &offset,
-          &srcMem, &offset, &dim, agent, direction, 1, &wait, active);
+          &srcMem, &offset, &dim, agent, direction, num_wait_events, wait_event, active);
       if (status != HSA_STATUS_SUCCESS) {
         LogPrintfError("DMA buffer failed with code %d", status);
         return false;
@@ -456,6 +458,8 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
       const hsa_signal_value_t kInitVal = size[2] * size[1];
       hsa_signal_t wait = gpu().Barriers().WaitSignal();
       hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitVal, gpu().timestamp());
+      uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1;
+      hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait;
 
       for (size_t z = 0; z < size[2]; ++z) {
         for (size_t y = 0; y < size[1]; ++y) {
@@ -466,7 +470,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
           hsa_status_t status = hsa_amd_memory_async_copy(
               (reinterpret_cast<address>(dst) + dstOffset), dstAgent,
               (reinterpret_cast<const_address>(src) + srcOffset), srcAgent,
-              size[0], 1, &wait, active);
+              size[0], num_wait_events, wait_event, active);
           gpu().setLastCommandSDMA(true) ;
           if (status != HSA_STATUS_SUCCESS) {
             LogPrintfError("DMA buffer failed with code %d", status);
@@ -640,8 +644,12 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
 
   hsa_signal_t wait = gpu().Barriers().WaitSignal();
   hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp());
+  uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1;
+  hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait;
+
   // Use SDMA to transfer the data
-  status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, size[0], 1, &wait, active);
+  status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent,
+      size[0], num_wait_events, wait_event, active);
   gpu().setLastCommandSDMA(true);
   // Explicit wait for now, until runtime could distinguish compute and sdma operations
   gpu().Barriers().WaitCurrent();

rocclr/device/rocm/rocdefs.hpp

@@ -33,6 +33,13 @@ static constexpr uint DeviceQueueMaskSize = 32;
 //! Set to match the number of pipes, which is 8.
 static constexpr uint kMaxAsyncQueues = 8;
 
+enum HwQueueEngine : uint32_t {
+  Compute   = 0,
+  SdmaRead  = 1,
+  SdmaWrite = 2,
+  Unknown   = 3
+};
+
 } // namespace roc
 
 #endif

rocclr/device/rocm/rocvirtual.cpp

@@ -871,7 +871,7 @@ void VirtualGPU::profilingBegin(amd::Command& command, bool drmProfiling) {
 */
 void VirtualGPU::profilingEnd(amd::Command& command) {
   if (command.profilingInfo().enabled_) {
-    if (timestamp_->getProfilingSignal() == nullptr) {
+    if (!timestamp_->HwProfiling()) {
       timestamp_->end();
     }
     command.setData(reinterpret_cast<void*>(timestamp_));
@@ -1186,6 +1186,8 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
   // Initialize signal for the barrier
   hsa_signal_t wait = Barriers().WaitSignal();
   hsa_signal_t active = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_);
+  uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1;
+  hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait;
 
   // Find the requested agent for the transfer
   hsa_agent_t agent = (cmd.cpu_access() ||
@@ -1194,7 +1196,7 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
 
   // Initiate a prefetch command
   hsa_status_t status = hsa_amd_svm_prefetch_async(
-      const_cast<void*>(cmd.dev_ptr()), cmd.count(), agent, 1, &wait, active);
+      const_cast<void*>(cmd.dev_ptr()), cmd.count(), agent, num_wait_events, wait_event, active);
 
   // Wait for the prefetch. Should skip wait, but may require extra tracking for kernel execution.
   if ((status != HSA_STATUS_SUCCESS) || !Barriers().WaitCurrent()) {
@@ -2134,9 +2136,6 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
         dim = i;
         iteration = sizes.global()[i] / 0xC0000000 + ((sizes.global()[i] % 0xC0000000) ? 1 : 0);
         globalStep = (sizes.global()[i] / sizes.local()[i]) / iteration * sizes.local()[dim];
-        if (timestamp_ != nullptr) {
-          timestamp_->setSplittedDispatch();
-        }
         break;
       }
     }

rocclr/device/rocm/rocvirtual.hpp

@@ -21,6 +21,7 @@
 #pragma once
 
 #include "platform/commandqueue.hpp"
+#include "rocdefs.hpp"
 #include "rocdevice.hpp"
 #include "utils/util.hpp"
 #include "hsa.h"
@@ -36,11 +37,15 @@ class Memory;
 class Timestamp;
 
 struct ProfilingSignal : public amd::HeapObject {
-  hsa_signal_t signal_; //!< HSA signal to track profiling information
-  Timestamp* ts_;       //!< Timestamp object associated with the signal
-  bool done_;           //!< True if signal is done
-
-  ProfilingSignal() : ts_(nullptr), done_(true) { signal_.handle = 0; }
+  hsa_signal_t  signal_;  //!< HSA signal to track profiling information
+  Timestamp*    ts_;      //!< Timestamp object associated with the signal
+  HwQueueEngine engine_;  //!< Engine used with this signal
+  bool          done_;    //!< True if signal is done
+  ProfilingSignal()
+    : ts_(nullptr)
+    , engine_(HwQueueEngine::Compute)
+    , done_(true)
+    { signal_.handle = 0; }
 };
 
 // Initial HSA signal value
@@ -67,13 +72,12 @@ inline bool WaitForSignal(hsa_signal_t signal) {
 // including EnqueueNDRangeKernel and clEnqueueCopyBuffer.
 class Timestamp {
  private:
-  uint64_t start_;
-  uint64_t end_;
-  ProfilingSignal* profilingSignal_;
-  hsa_agent_t agent_;
   static double ticksToTime_;
-  bool splittedDispatch_;
-  std::vector<hsa_signal_t> splittedSignals_;
+
+  uint64_t    start_;
+  uint64_t    end_;
+  hsa_agent_t agent_;
+  std::vector<ProfilingSignal*> signals_;
 
  public:
   uint64_t getStart() {
@@ -86,21 +90,15 @@ class Timestamp {
     return end_;
   }
 
-  void setProfilingSignal(ProfilingSignal* signal) {
-    profilingSignal_ = signal;
-    if (splittedDispatch_) {
-      splittedSignals_.push_back(profilingSignal_->signal_);
-    }
-  }
-  const ProfilingSignal* getProfilingSignal() const { return profilingSignal_; }
+  void AddProfilingSignal(ProfilingSignal* signal) { signals_.push_back(signal); }
+
+  const bool HwProfiling() const { return (signals_.size() > 0) ? true : false; }
 
   void setAgent(hsa_agent_t agent) { agent_ = agent; }
 
   Timestamp()
-    : start_(0)
-    , end_(0)
-    , profilingSignal_(nullptr)
-    , splittedDispatch_(false) {
+    : start_(std::numeric_limits<uint64_t>::max())
+    , end_(0) {
     agent_.handle = 0;
   }
 
@@ -108,51 +106,30 @@ class Timestamp {
 
   //! Finds execution ticks on GPU
   void checkGpuTime() {
-    if (profilingSignal_ != nullptr) {
-      hsa_amd_profiling_dispatch_time_t time;
+    if (HwProfiling()) {
+      hsa_amd_profiling_dispatch_time_t time = {};
 
-      if (splittedDispatch_) {
-        uint64_t start = std::numeric_limits<uint64_t>::max();
-        uint64_t end = 0;
-        for (auto it = splittedSignals_.begin(); it < splittedSignals_.end(); it++) {
-          if (hsa_signal_load_relaxed(profilingSignal_->signal_) > 0) {
-            WaitForSignal(*it);
-          }
-          hsa_amd_profiling_get_dispatch_time(agent_, *it, &time);
-          if ((time.end - time.start) == 0) {
-            hsa_amd_profiling_async_copy_time_t time_sdma = {};
-            hsa_amd_profiling_get_async_copy_time(profilingSignal_->signal_, &time_sdma);
-            time.start = time_sdma.start;
-            time.end = time_sdma.end;
-          }
-          if (time.start < start) {
-            start = time.start;
-          }
-          if (time.end > end) {
-            end = time.end;
-          }
+      uint64_t start = std::numeric_limits<uint64_t>::max();
+      uint64_t end = 0;
+      for (auto it : signals_) {
+        if (hsa_signal_load_relaxed(it->signal_) > 0) {
+          WaitForSignal(it->signal_);
         }
-        start_ = start * ticksToTime_;
-        end_ = end * ticksToTime_;
-      } else {
-        // If the signalValue is the same as initial set value, it means its not written to
-        if (hsa_signal_load_relaxed(profilingSignal_->signal_) > 0) {
-          WaitForSignal(profilingSignal_->signal_);
-        }
-        hsa_amd_profiling_get_dispatch_time(agent_, profilingSignal_->signal_, &time);
+        hsa_amd_profiling_get_dispatch_time(agent_, it->signal_, &time);
         if ((time.end - time.start) == 0) {
           hsa_amd_profiling_async_copy_time_t time_sdma = {};
-          hsa_amd_profiling_get_async_copy_time(profilingSignal_->signal_, &time_sdma);
-          start_ = time_sdma.start * ticksToTime_;
-          end_ = time_sdma.end * ticksToTime_;
-        } else {
-          start_ = time.start * ticksToTime_;
-          end_ = time.end * ticksToTime_;
+          hsa_amd_profiling_get_async_copy_time(it->signal_, &time_sdma);
+          time.start = time_sdma.start;
+          time.end = time_sdma.end;
         }
+        start = std::min(time.start, start);
+        end = std::max(time.end, end);
+        it->ts_ = nullptr;
+        it->done_ = true;
       }
-      profilingSignal_->ts_ = nullptr;
-      profilingSignal_->done_ = true;
-      profilingSignal_ = nullptr;
+      signals_.clear();
+      start_ = start * ticksToTime_;
+      end_ = end * ticksToTime_;
     }
   }
 
@@ -162,9 +139,6 @@ class Timestamp {
   // End a timestamp (get timestamp from OS)
   void end() { end_ = amd::Os::timeNanos(); }
 
-  bool isSplittedDispatch() const { return splittedDispatch_; }
-  void setSplittedDispatch() { splittedDispatch_ = true; }
-
   static void setGpuTicksToTime(double ticksToTime) { ticksToTime_ = ticksToTime; }
   static double getGpuTicksToTime() { return ticksToTime_; }
 };
@@ -268,7 +242,7 @@ class VirtualGPU : public device::VirtualDevice {
           sdma_profiling_ = true;
         }
         signal_list_[current_id_]->ts_ = ts;
-        ts->setProfilingSignal(signal_list_[current_id_]);
+        ts->AddProfilingSignal(signal_list_[current_id_]);
         ts->setAgent(agent_);
       }
       return signal_list_[current_id_]->signal_;
@@ -278,7 +252,11 @@ class VirtualGPU : public device::VirtualDevice {
     bool WaitCurrent() { return WaitIndex(current_id_); }
 
     //! Returns the last submitted signal for a wait
-    hsa_signal_t WaitSignal() const { return signal_list_[current_id_]->signal_; }
+    hsa_signal_t WaitSignal() {
+      //! @note Currently wait on CPU unconditionally to avoid a negative performance impact
+      WaitCurrent();
+      return hsa_signal_t{};
+    }
 
    private:
     //! Wait for the next active signal