rocr: Add thread priority for AsyncEventHandler

Set priority to maximum for signal event handler and minimum for
exceptions event handler.

Change-Id: I1b982d3c2e4c880fafc073fe1a542d01692a6fdc

runtime/hsa-runtime/core/runtime/runtime.cpp

@@ -809,14 +809,17 @@ hsa_status_t Runtime::SetAsyncSignalHandler(hsa_signal_t signal,
                                             void* arg) {
 
   struct AsyncEventsInfo* asyncInfo = &asyncSignals_;
+  int priority = runtime_singleton_->flag().async_events_thread_priority();
 
   if (signal.handle != 0) {
     // Indicate that this signal is in use.
     hsa_signal_handle(signal)->Retain();
 
     core::Signal* coreSignal = core::Signal::Convert(signal);
-    if (coreSignal->EopEvent() && coreSignal->EopEvent()->EventData.EventType != HSA_EVENTTYPE_SIGNAL)
+    if (coreSignal->EopEvent() && coreSignal->EopEvent()->EventData.EventType != HSA_EVENTTYPE_SIGNAL) {
+      priority = os::OS_THREAD_PRIORITY_DEFAULT;
       asyncInfo = &asyncExceptions_;
+    }
   }
 
   ScopedAcquire<HybridMutex> scope_lock(&asyncInfo->control.lock);
@@ -835,7 +838,7 @@ hsa_status_t Runtime::SetAsyncSignalHandler(hsa_signal_t signal,
     // Start event monitoring thread
     asyncInfo->control.exit = false;
     asyncInfo->control.async_events_thread_ =
-        os::CreateThread(AsyncEventsLoop, asyncInfo);
+        os::CreateThread(AsyncEventsLoop, asyncInfo, 0, priority);
     if (asyncInfo->control.async_events_thread_ == NULL) {
       assert(false && "Asyncronous events thread creation error.");
       return HSA_STATUS_ERROR_OUT_OF_RESOURCES;

runtime/hsa-runtime/core/util/flag.h

@@ -257,6 +257,22 @@ class Flag {
     /* hsa_signal_wait_relaxed abort timeout  */
     var = os::GetEnvVar("HSA_SIGNAL_WAIT_ABORT_TIMEOUT");
     signal_abort_timeout_ = var.empty() ? 0 : atoi(var.c_str());
+
+    /* Valid inputs are 0-99, HIGH, MAX */
+    var = os::GetEnvVar("HSA_ASYNCEVENTS_THREAD_PRIORITY");
+    async_events_thread_priority_ = os::OS_THREAD_PRIORITY_DEFAULT;
+    if (var == "MAX") {
+      async_events_thread_priority_ = os::OS_THREAD_PRIORITY_MAX;
+    } else if (var == "HIGH") {
+      async_events_thread_priority_ = os::OS_THREAD_PRIORITY_HIGH;
+    } else if (var != "") {
+      char* end;
+      int input = strtol(var.c_str(), &end, 10);
+      if (input >= 0 && input <= 99)
+        async_events_thread_priority_ = input;
+      else
+        fprintf(stderr, "Failed to parse HSA_ASYNCEVENTS_THREAD_PRIORITY");
+    }
   }
 
   void parse_masks(uint32_t maxGpu, uint32_t maxCU) {
@@ -373,6 +389,8 @@ class Flag {
 
   uint32_t signal_abort_timeout() const { return signal_abort_timeout_; }
 
+  int async_events_thread_priority() const { return async_events_thread_priority_; }
+
  private:
   bool check_flat_scratch_;
   bool enable_vm_fault_message_;
@@ -404,6 +422,7 @@ class Flag {
   bool wait_any_;
   bool dev_mem_queue_;
   uint32_t signal_abort_timeout_;
+  int  async_events_thread_priority_;
 
   SDMA_OVERRIDE enable_sdma_;
   SDMA_OVERRIDE enable_peer_sdma_;

runtime/hsa-runtime/core/util/lnx/os_linux.cpp

@@ -90,7 +90,10 @@ void* __stdcall ThreadTrampoline(void* arg) {
 // Thread container allows multiple waits and separate close (destroy).
 class os_thread {
  public:
-  explicit os_thread(ThreadEntry function, void* threadArgument, uint stackSize)
+  explicit os_thread(ThreadEntry function,
+                      void* threadArgument,
+                      uint stackSize,
+                      int priority)
       : thread(0), lock(nullptr), state(RUNNING) {
     int err;
     lock = CreateMutex();
@@ -162,6 +165,38 @@ class os_thread {
         return;
       }
     } while (stackSize < 20 * 1024 * 1024);
+
+    struct sched_param param = {};
+    if (priority != OS_THREAD_PRIORITY_DEFAULT) {
+      int set_priority;
+      int max_priority = sched_get_priority_max(SCHED_FIFO);
+
+      if (priority == OS_THREAD_PRIORITY_MAX)
+        set_priority = max_priority;
+      else if (priority == OS_THREAD_PRIORITY_HIGH)
+        set_priority = max_priority - 1;
+      else if (priority > max_priority)
+        set_priority = max_priority;
+      else
+        set_priority = priority;
+
+      param.sched_priority = set_priority;
+      if (pthread_setschedparam(thread, SCHED_FIFO, &param)) {
+        fprintf(stderr, "pthread_setschedparam failed\n");
+        return;
+      }
+
+      int policy = 0;
+      if (pthread_getschedparam(thread, &policy, &param))
+        fprintf(stderr, "pthread_getschedparam failed: %s\n", strerror(err));
+
+      if (policy != SCHED_FIFO || param.sched_priority != set_priority)
+        fprintf(stderr, "Failed to adjust thread priority (policy:%s requested:%d current:%d)\n",
+                          policy == SCHED_FIFO ? "FIFO" :
+                          policy == SCHED_OTHER ? "OTHER" :
+                          policy == SCHED_RR ? "RR" : "Unknown",
+                          set_priority, param.sched_priority);
+    }
   }
 
   os_thread(os_thread&& rhs) {
@@ -413,8 +448,8 @@ void uSleep(int delayInUs) { usleep(delayInUs); }
 
 void YieldThread() { sched_yield(); }
 
-Thread CreateThread(ThreadEntry function, void* threadArgument, uint stackSize) {
-  os_thread* result = new os_thread(function, threadArgument, stackSize);
+Thread CreateThread(ThreadEntry function, void* threadArgument, uint stackSize, int priority) {
+  os_thread* result = new os_thread(function, threadArgument, stackSize, priority);
   if (!result->Valid()) {
     delete result;
     return nullptr;

runtime/hsa-runtime/core/util/os.h

@@ -58,6 +58,12 @@ typedef void* SharedMutex;
 typedef void* Thread;
 typedef void* EventHandle;
 
+typedef enum {
+  OS_THREAD_PRIORITY_DEFAULT    = -1,
+  OS_THREAD_PRIORITY_HIGH       = 254,
+  OS_THREAD_PRIORITY_MAX        = 255,
+} ThreadPriority;
+
 enum class os_t { OS_WIN = 0, OS_LINUX, COUNT };
 static __forceinline std::underlying_type<os_t>::type os_index(os_t val) {
   return std::underlying_type<os_t>::type(val);
@@ -209,9 +215,10 @@ typedef void (*ThreadEntry)(void*);
 /// @param: entry_argument(Input), a pointer to the argument of the thread
 /// function.
 /// @param: stack_size(Input), size of the thread's stack, 0 by default.
+/// @param: priority(Input), thread priority.
 /// @return: Thread, a handle to thread created.
 Thread CreateThread(ThreadEntry entry_function, void* entry_argument,
-                    uint stack_size = 0);
+                    uint stack_size = 0, int priority = OS_THREAD_PRIORITY_DEFAULT);
 
 /// @brief: Destroys the thread.
 /// @param: thread(Input), thread handle to what will be destroyed.

runtime/hsa-runtime/core/util/win/os_win.cpp

@@ -151,7 +151,7 @@ unsigned __stdcall ThreadTrampoline(void* arg) {
 }
 
 Thread CreateThread(ThreadEntry entry_function, void* entry_argument,
-                    uint stack_size) {
+                    uint stack_size, int priority_unused) {
   ThreadArgs* thread_args = new ThreadArgs();
   thread_args->entry_args = entry_argument;
   thread_args->entry_function = entry_function;