From 95ee8b36d7ecc6103fb35ac642822225c968c464 Mon Sep 17 00:00:00 2001 From: David Yat Sin Date: Wed, 11 Dec 2024 17:51:21 +0000 Subject: [PATCH] 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: 5da1889fb7bacf064c56a7b8af9374e5f3fb2ad2] --- .../runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp | 5 +++++ .../runtime/hsa-runtime/core/runtime/default_signal.cpp | 3 --- .../runtime/hsa-runtime/core/runtime/runtime.cpp | 2 +- 3 files changed, 6 insertions(+), 4 deletions(-) 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],