rocr: Avoid deadlock due to queue signal not updated

Make sure waiting_ count for queue signal is always > 0 so that we
always call hsaKmtWaitOnEvent to force hsaKmtWaitOnEvent to return.

Remove incorrect warning print when running in debug mode.

Call internal Signal::WaitAny instead of AMD::hsa_amd_signal_wait_any
to avoid extra function calls.

Change-Id: I9e41b704643e4e8ee7402b1379b1c30ff4c544ef


[ROCm/ROCR-Runtime commit: 5da1889fb7]

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

@@ -270,6 +270,10 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
     assert(exception_signal_ != nullptr && "Should have thrown!\n");
   }
 
+  // Make sure the queue signal always has a waiting_ > 0 so that
+  // so that we call hsakmtSetEvent to force hsaKmtWaitOnEvent to return.
+  exception_signal_->WaitingInc();
+
   // Ensure the amd_queue_ is fully initialized before creating the KFD queue.
   // This ensures that the debugger can access the fields once it detects there
   // is a KFD queue. The debugger may access the aperture addresses, queue
@@ -374,6 +378,7 @@ AqlQueue::~AqlQueue() {
   if (queue_scratch_.alt_queue_base) agent_->ReleaseQueueAltScratch(queue_scratch_);
 
   FreeRegisteredRingBuffer();
+  exception_signal_->WaitingDec();
   exception_signal_->DestroySignal();
   HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal);
   if (core::g_use_interrupt_wait) {

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

@@ -89,9 +89,6 @@ hsa_signal_value_t BusyWaitSignal::WaitRelaxed(hsa_signal_condition_t condition,
   const uint32_t &signal_abort_timeout =
     core::Runtime::runtime_singleton_->flag().signal_abort_timeout();
 
-  debug_warning_n((!g_use_interrupt_wait || isIPC()) &&
-                  "Use of non-host signal in host signal wait API.", 10);
-
   timer::fast_clock::time_point start_time, time;
   start_time = timer::fast_clock::now();
 

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

@@ -1636,7 +1636,7 @@ void Runtime::AsyncEventsLoop(void* _eventsInfo) {
                           &value);
     } else {
      if (core::Runtime::runtime_singleton_->flag().wait_any()) {
-      index = AMD::hsa_amd_signal_wait_any(
+      index = Signal::WaitAny(
                           uint32_t(async_events_.Size()),
                           &async_events_.signal_[0],
                           &async_events_.cond_[0],