From 5288e46340a498e63ac451c310467387972bcf59 Mon Sep 17 00:00:00 2001
From: Sean Keely <Sean.Keely@amd.com>
Date: Wed, 22 Mar 2017 14:32:13 -0500
Subject: [PATCH] Refactor signal_wait timing code and respect small timeouts.

Optimized for Gromacs and SHOC.

Change-Id: Ib674710268b41003259711a0e42d3e770a82018d


[ROCm/ROCR-Runtime commit: c4544906b93d0be14164e0a25679b2251a61ec17]
---
 .../core/runtime/interrupt_signal.cpp         | 37 ++++++++++---------
 .../hsa-runtime/core/runtime/signal.cpp       | 35 +++++++++---------
 2 files changed, 37 insertions(+), 35 deletions(-)

diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp
index 949ed8af25..84d20cf28d 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp
+++ b/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);
   }
 }
 
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/signal.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/signal.cpp
index 222d7e5b06..8b26396849 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/signal.cpp
+++ b/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);
   }
 }