SWDEV-358915 - Wait for 10 micro secs, if the next engine is SDMA.

Change-Id: Ib3ed91510e62090f47d4f55be8de635caaee083d


[ROCm/clr commit: 61b6a94627]

projects/clr/rocclr/device/rocm/rocvirtual.cpp

@@ -475,9 +475,13 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(
 // ================================================================================================
 std::vector<hsa_signal_t>& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngine engine) {
   bool explicit_wait = false;
+  bool sdma_wait = false;
   // Reset all current waiting signals
   waiting_signals_.clear();
 
+  if(engine != HwQueueEngine::Compute)
+    sdma_wait = true;
+
   // Does runtime switch the active engine?
   if (engine != engine_) {
     // Yes, return the signal from the previous operation for a wait
@@ -518,7 +522,7 @@ std::vector<hsa_signal_t>& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi
       if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) {
         const Settings& settings = gpu_.dev().settings();
         // Actively wait on CPU to avoid extra overheads of signal tracking on GPU
-        if (!WaitForSignal<true>(external_signals_[i]->signal_)) {
+        if (!WaitForSignal<true>(external_signals_[i]->signal_, false, sdma_wait)) {
           if (settings.cpu_wait_for_signal_) {
             // Wait on CPU for completion if requested
             CpuWaitForSignal(external_signals_[i]);

projects/clr/rocclr/device/rocm/rocvirtual.hpp

@@ -44,15 +44,20 @@ constexpr static hsa_signal_value_t kInitSignalValueOne = 1;
 constexpr static uint64_t kTimeout100us = 100 * K;
 constexpr static uint64_t kUnlimitedWait = std::numeric_limits<uint64_t>::max();
 
+// Active wait time out incase same sdma engine is used again,
+// then just wait instead of adding dependency wait signal.
+constexpr static uint64_t kSDMAEngineTimeout = 10;
+
 template <bool active_wait_timeout = false>
-inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false) {
+inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool sdma_wait = false) {
   if (hsa_signal_load_relaxed(signal) > 0) {
     uint64_t timeout = kTimeout100us;
     if (active_wait) {
       timeout = kUnlimitedWait;
     }
     if (active_wait_timeout) {
-      timeout = ROC_ACTIVE_WAIT_TIMEOUT * K;
+      // If diff engine, wait to 10 ms. Otherwise no wait
+      timeout = (sdma_wait ? kSDMAEngineTimeout : ROC_ACTIVE_WAIT_TIMEOUT) * K;
       if (timeout == 0) {
         return false;
       }