Refactor signal_wait timing code and respect small timeouts.

Optimized for Gromacs and SHOC.

Change-Id: Ib674710268b41003259711a0e42d3e770a82018d


[ROCm/ROCR-Runtime commit: c4544906b9]

projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp

@@ -171,25 +171,26 @@ hsa_signal_value_t InterruptSignal::WaitRelaxed(
     if (condition_met) return hsa_signal_value_t(value);
 
     timer::fast_clock::time_point time = timer::fast_clock::now();
-    if (time - start_time > kMaxElapsed) {
-      if (time - start_time > fast_timeout) {
-        value = atomic::Load(&signal_.value, std::memory_order_relaxed);
-        return hsa_signal_value_t(value);
-      }
-      if (wait_hint != HSA_WAIT_STATE_ACTIVE) {
-        uint32_t wait_ms;
-        auto time_remaining = fast_timeout - (time - start_time);
-        if ((timeout == -1) ||
-            (time_remaining > std::chrono::milliseconds(uint32_t(-1))))
-          wait_ms = uint32_t(-1);
-        else
-          wait_ms = timer::duration_cast<std::chrono::milliseconds>(
-                        time_remaining).count();
-        hsaKmtWaitOnEvent(event_, wait_ms);
-      }
-    } else {
-      os::uSleep(20);
+    if (time - start_time > fast_timeout) {
+      value = atomic::Load(&signal_.value, std::memory_order_relaxed);
+      return hsa_signal_value_t(value);
     }
+    
+    if (wait_hint == HSA_WAIT_STATE_ACTIVE) {
+      continue;
+    }
+
+    if (time - start_time < kMaxElapsed) {
+    //  os::uSleep(20);
+      continue;
+    }
+
+    uint32_t wait_ms;
+    auto time_remaining = fast_timeout - (time - start_time);
+    uint64_t ct=timer::duration_cast<std::chrono::milliseconds>(
+      time_remaining).count();
+    wait_ms = (ct>0xFFFFFFFEu) ? 0xFFFFFFFEu : ct;
+    hsaKmtWaitOnEvent(event_, wait_ms);
   }
 }
 

projects/rocr-runtime/runtime/hsa-runtime/core/runtime/signal.cpp

@@ -160,24 +160,25 @@ uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals,
     }
 
     timer::fast_clock::time_point time = timer::fast_clock::now();
-    if (time - start_time > kMaxElapsed) {
-      if (time - start_time > fast_timeout) {
-        return uint32_t(-1);
-      }
-      if (wait_hint != HSA_WAIT_STATE_ACTIVE) {
-        uint32_t wait_ms;
-        auto time_remaining = fast_timeout - (time - start_time);
-        if ((timeout == -1) ||
-            (time_remaining > std::chrono::milliseconds(uint32_t(-1))))
-          wait_ms = uint32_t(-1);
-        else
-          wait_ms = timer::duration_cast<std::chrono::milliseconds>(
-                        time_remaining).count();
-        hsaKmtWaitOnMultipleEvents(evts, unique_evts, false, wait_ms);
-      }
-    } else {
-      os::uSleep(20);
+    if (time - start_time > fast_timeout) {
+      return uint32_t(-1);
     }
+
+    if (wait_hint == HSA_WAIT_STATE_ACTIVE) {
+      continue;
+    }
+
+    if (time - start_time < kMaxElapsed) {
+    //  os::uSleep(20);
+      continue;
+    }
+
+    uint32_t wait_ms;
+    auto time_remaining = fast_timeout - (time - start_time);
+    uint64_t ct=timer::duration_cast<std::chrono::milliseconds>(
+      time_remaining).count();
+    wait_ms = (ct>0xFFFFFFFEu) ? 0xFFFFFFFEu : ct;
+    hsaKmtWaitOnMultipleEvents(evts, unique_evts, false, wait_ms);
   }
 }