SWDEV-290384 - Enable active wait on CPU if HIP requested

Change-Id: Idea5adf7a4705cb999da6785e6229fe3200dce17

rocclr/device/rocm/rocdevice.cpp

@@ -2546,7 +2546,8 @@ bool Device::IsHwEventReady(const amd::Event& event, bool wait) const {
     ClPrint(amd::LOG_INFO, amd::LOG_SIG, "No HW event");
     return false;
   } else if (wait) {
-    WaitForSignal(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_);
+    auto* vdev = event.command().queue()->vdev();
+    WaitForSignal(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, vdev->ActiveWait());
     return true;
   }
   return (hsa_signal_load_relaxed(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_) <= 0);

rocclr/device/rocm/rocvirtual.cpp

@@ -508,7 +508,7 @@ bool VirtualGPU::HwQueueTracker::CpuWaitForSignal(ProfilingSignal* signal) {
     } else {
       ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "[%zx]!\t Host wait on completion_signal=0x%zx",
               std::this_thread::get_id(), signal->signal_.handle);
-      if (!WaitForSignal(signal->signal_)) {
+      if (!WaitForSignal(signal->signal_, gpu_.ActiveWait())) {
         LogPrintfError("Failed signal [0x%lx] wait", signal->signal_);
         return false;
       }

rocclr/device/rocm/rocvirtual.hpp

@@ -45,34 +45,34 @@ constexpr static uint64_t kTimeout100us = 100 * K;
 constexpr static uint64_t kUnlimitedWait = std::numeric_limits<uint64_t>::max();
 
 template <bool active_wait_timeout = false>
-inline bool WaitForSignal(hsa_signal_t signal) {
+inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false) {
   if (hsa_signal_load_relaxed(signal) > 0) {
-    if (active_wait_timeout) {
-      uint64_t timeout = ROC_ACTIVE_WAIT_TIMEOUT * K;
+    uint64_t timeout = kTimeout100us;
+    if (active_wait) {
+      timeout = kUnlimitedWait;
+    } else if (active_wait_timeout) {
+      timeout = ROC_ACTIVE_WAIT_TIMEOUT * K;
       if (timeout == 0) {
         return false;
       }
-      ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host active wait for Signal = (0x%lx) for %d us",
-              signal.handle, ROC_ACTIVE_WAIT_TIMEOUT);
+    }
 
-      if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
-                                    timeout, HSA_WAIT_STATE_ACTIVE) != 0) {
+    ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host active wait for Signal = (0x%lx) for %d ns",
+            signal.handle, timeout);
+
+    // Active wait with a timeout
+    if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
+                                  timeout, HSA_WAIT_STATE_ACTIVE) != 0) {
+      if (active_wait_timeout) {
         return false;
       }
-    } else {
-
-      uint64_t timeout = kTimeout100us;
-      ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host wait until Signal = (0x%lx) decremented",
+      ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host blocked wait for Signal = (0x%lx)",
               signal.handle);
 
-      // Active wait with a timeout
+      // Wait until the completion with CPU suspend
       if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
-                                    timeout, HSA_WAIT_STATE_ACTIVE) != 0) {
-        // Wait until the completion with CPU suspend
-        if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
-                                      kUnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) {
-          return false;
-        }
+                                    kUnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) {
+        return false;
       }
     }
   }