From c02e8c67df450c09655b79b8098bacb314c1bbd9 Mon Sep 17 00:00:00 2001
From: Sean Keely <Sean.Keely@amd.com>
Date: Wed, 8 Mar 2017 02:59:39 -0600
Subject: [PATCH] Adjust signal sleep to reflect null kernel latency. 
 Performance tested on Gromacs.

Change-Id: I3851148ee8544b15d840f2c26ca73a83f8d0df2e


[ROCm/ROCR-Runtime commit: 426d41e27c80b8482996e380d090369e0849b7f4]
---
 .../runtime/hsa-runtime/core/runtime/default_signal.cpp    | 1 +
 .../runtime/hsa-runtime/core/runtime/interrupt_signal.cpp  | 7 ++++---
 .../runtime/hsa-runtime/core/runtime/signal.cpp            | 6 +++---
 .../runtime/hsa-runtime/core/util/lnx/os_linux.cpp         | 2 ++
 projects/rocr-runtime/runtime/hsa-runtime/core/util/os.h   | 5 +++++
 .../runtime/hsa-runtime/core/util/win/os_win.cpp           | 2 ++
 6 files changed, 17 insertions(+), 6 deletions(-)

diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp
index 9b81de360b..097d204356 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/default_signal.cpp
@@ -131,6 +131,7 @@ hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition,
       value = atomic::Load(&signal_.value, std::memory_order_relaxed);
       return hsa_signal_value_t(value);
     }
+    os::uSleep(20);
   }
 }
 
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 eb07bcc533..949ed8af25 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
@@ -133,9 +133,8 @@ hsa_signal_value_t InterruptSignal::WaitRelaxed(
   timer::fast_clock::time_point start_time = timer::fast_clock::now();
 
   // Set a polling timeout value
-  // Exact time is not hugely important, it should just be a short while which
-  // is smaller than the thread scheduling quantum (usually around 16ms)
-  const timer::fast_clock::duration kMaxElapsed = std::chrono::milliseconds(5);
+  // Should be a few times bigger than null kernel latency
+  const timer::fast_clock::duration kMaxElapsed = std::chrono::microseconds(200);
 
   uint64_t hsa_freq;
   HSA::hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &hsa_freq);
@@ -188,6 +187,8 @@ hsa_signal_value_t InterruptSignal::WaitRelaxed(
                         time_remaining).count();
         hsaKmtWaitOnEvent(event_, wait_ms);
       }
+    } else {
+      os::uSleep(20);
     }
   }
 }
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 0874080da3..222d7e5b06 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/signal.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/signal.cpp
@@ -101,9 +101,7 @@ uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals,
   timer::fast_clock::time_point start_time = timer::fast_clock::now();
 
   // Set a polling timeout value
-  // Exact time is not hugely important, it should just be a short while which
-  // is smaller than the thread scheduling quantum (usually around 16ms)
-  const timer::fast_clock::duration kMaxElapsed = std::chrono::milliseconds(5);
+  const timer::fast_clock::duration kMaxElapsed = std::chrono::microseconds(200);
 
   // Convert timeout value into the fast_clock domain
   uint64_t hsa_freq;
@@ -177,6 +175,8 @@ uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals,
                         time_remaining).count();
         hsaKmtWaitOnMultipleEvents(evts, unique_evts, false, wait_ms);
       }
+    } else {
+      os::uSleep(20);
     }
   }
 }
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/util/lnx/os_linux.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/util/lnx/os_linux.cpp
index 88f339a93e..3200cf2cdf 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/util/lnx/os_linux.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/util/lnx/os_linux.cpp
@@ -115,6 +115,8 @@ void DestroyMutex(Mutex lock) {
 
 void Sleep(int delay_in_millisec) { usleep(delay_in_millisec * 1000); }
 
+void uSleep(int delayInUs) { usleep(delayInUs); }
+
 void YieldThread() { sched_yield(); }
 
 struct ThreadArgs {
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/util/os.h b/projects/rocr-runtime/runtime/hsa-runtime/core/util/os.h
index 279468b018..4421b4a3e5 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/util/os.h
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/util/os.h
@@ -114,6 +114,11 @@ void DestroyMutex(Mutex lock);
 /// @return: void.
 void Sleep(int delayInMs);
 
+/// @brief: Puts current thread to sleep.
+/// @param: delayInMs(Input), time in millisecond for sleeping.
+/// @return: void.
+void uSleep(int delayInUs);
+
 /// @brief: Yields current thread.
 /// @param: void.
 /// @return: void.
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/util/win/os_win.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/util/win/os_win.cpp
index d97bff0ce8..a55b3af49a 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/util/win/os_win.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/util/win/os_win.cpp
@@ -98,6 +98,8 @@ void DestroyMutex(Mutex lock) { CloseHandle(*(::HANDLE*)&lock); }
 
 void Sleep(int delay_in_millisecond) { ::Sleep(delay_in_millisecond); }
 
+void uSleep(int delayInUs) { ::Sleep(delayInUs / 1000); }
+
 void YieldThread() { ::Sleep(0); }
 
 struct ThreadArgs {