From 2babcfbdbbed2935adeda465bc666cf35d58d4c2 Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Mon, 19 Jul 2021 15:05:24 -0400
Subject: [PATCH] SWDEV-290384 - Enable active wait on CPU if HIP requested

Change-Id: Idea5adf7a4705cb999da6785e6229fe3200dce17
---
 rocclr/device/rocm/rocdevice.cpp  |  3 ++-
 rocclr/device/rocm/rocvirtual.cpp |  2 +-
 rocclr/device/rocm/rocvirtual.hpp | 36 +++++++++++++++----------------
 3 files changed, 21 insertions(+), 20 deletions(-)

diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp
index 69ab3590df..c72e14c285 100644
--- a/rocclr/device/rocm/rocdevice.cpp
+++ b/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);
diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp
index 1c0773bf50..9cde1f30a2 100644
--- a/rocclr/device/rocm/rocvirtual.cpp
+++ b/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;
       }
diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp
index f5c8666a93..0d5b410105 100644
--- a/rocclr/device/rocm/rocvirtual.hpp
+++ b/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;
       }
     }
   }