diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
index 455d182d1b..744d8bf072 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+++ b/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) {
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 5881f0e03f..827acc6e9f 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
@@ -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();
 
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp
index e3d4794ca5..f86169c542 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp
+++ b/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],