Adjust signal sleep to reflect null kernel latency. Performance tested on Gromacs.

Change-Id: I3851148ee8544b15d840f2c26ca73a83f8d0df2e


[ROCm/ROCR-Runtime commit: 426d41e27c]

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);
   }
 }
 

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);
     }
   }
 }

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);
     }
   }
 }

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 {

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.

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 {