SWDEV-247372 - Active wait timeout env var

- Create an env var ROC_ACTIVE_WAIT_TIMEOUT to set active wait timeout
- Record profiling informaion if marker_ts_ property is valid.

Change-Id: If0d8aec8d9b0715027cf0f7c3dc8a4c722a6bae6

rocclr/device/rocm/rocvirtual.cpp

@@ -122,7 +122,7 @@ void Timestamp::checkGpuTime() {
       }
       // Avoid profiling data for the sync barrier, in tiny performance tests the first call
       // to ROCr is very slow and that also affects the overall performance of the callback thread
-      if (command().GetBatchHead() == nullptr) {
+      if (command().GetBatchHead() == nullptr || command().profilingInfo().marker_ts_) {
         hsa_amd_profiling_dispatch_time_t time = {};
         if (it->engine_ == HwQueueEngine::Compute) {
           hsa_amd_profiling_get_dispatch_time(gpu()->gpu_device(), it->signal_, &time);
@@ -132,8 +132,11 @@ void Timestamp::checkGpuTime() {
           time.start = time_sdma.start;
           time.end = time_sdma.end;
         }
+
         start = std::min(time.start, start);
         end = std::max(time.end, end);
+        ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Signal = (0x%lx), start = %ld, "
+          "end = %ld", it->signal_.handle, start, end);
       }
       it->ts_ = nullptr;
       it->done_ = true;
@@ -452,8 +455,8 @@ std::vector<hsa_signal_t>& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi
       // Early signal status check
       if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) {
         const Settings& settings = gpu_.dev().settings();
-        // Actively wait on CPU for 750 us to avoid extra overheads of signal tracking on GPU
-        if (!WaitForSignal<kTimeout750us>(external_signals_[i]->signal_)) {
+        // Actively wait on CPU to avoid extra overheads of signal tracking on GPU
+        if (!WaitForSignal<true>(external_signals_[i]->signal_)) {
           if (settings.cpu_wait_for_signal_) {
             // Wait on CPU for completion if requested
             CpuWaitForSignal(external_signals_[i]);

rocclr/device/rocm/rocvirtual.hpp

@@ -52,30 +52,42 @@ struct ProfilingSignal : public amd::HeapObject {
 constexpr static hsa_signal_value_t kInitSignalValueOne = 1;
 
 // Timeouts for HSA signal wait
-constexpr static uint64_t kTimeout100us = 100000;
-constexpr static uint64_t kTimeout750us = 750000;
+constexpr static uint64_t kTimeout100us = 100 * K;
 constexpr static uint64_t kUnlimitedWait = std::numeric_limits<uint64_t>::max();
 
-template <uint64_t wait_time = 0>
+template <bool active_wait_timeout = false>
 inline bool WaitForSignal(hsa_signal_t signal) {
-  if (wait_time != 0) {
-    if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
-                                  wait_time, HSA_WAIT_STATE_ACTIVE) != 0) {
-      return false;
-    }
-  } else {
-    uint64_t timeout = (ROC_ACTIVE_WAIT) ? kUnlimitedWait : kTimeout100us;
-
-    // Active wait with a timeout
-    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) {
+  if (hsa_signal_load_relaxed(signal) > 0) {
+    if (active_wait_timeout) {
+      uint64_t 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) {
+        return false;
+      }
+    } else {
+
+      uint64_t timeout = kTimeout100us;
+      ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host wait until Signal = (0x%lx) decremented",
+              signal.handle);
+
+      // Active wait with a timeout
+      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;
+        }
+      }
     }
   }
+
   return true;
 }
 

rocclr/utils/flags.hpp

@@ -233,8 +233,8 @@ release(uint, HIP_HIDDEN_FREE_MEM, 0,                                         \
         "0 = Disable")                                                        \
 release(size_t, GPU_FORCE_BLIT_COPY_SIZE, 0,                                  \
         "Size in KB of the threshold below which to force blit instead for sdma") \
-release(bool, ROC_ACTIVE_WAIT, false,                                         \
-        "Forces unconditional active wait for GPU")                           \
+release(uint, ROC_ACTIVE_WAIT_TIMEOUT, 750,                                   \
+        "Forces active wait of GPU interrup for the timeout(us)")             \
 release(bool, ROC_ENABLE_LARGE_BAR, true,                                     \
         "Enable Large Bar if supported by the device")                        \
 release(bool, ROC_CPU_WAIT_FOR_SIGNAL, true,                                  \