From e06dd39d89c93c463cc093f5a4c656517b98ca66 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Wed, 25 Aug 2021 22:13:41 -0500 Subject: [PATCH] Correct queue exception hander termination. Set handler state to terminated before exiting. Also simplify scratch handler exit loop. Change-Id: I0a80c8a1899e8b60a6e7aa6989ba28de42ba31e7 --- runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 8e0e602e13..48c7d23d6a 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -330,11 +330,10 @@ AqlQueue::~AqlQueue() { // Remove error handler synchronously. // Sequences error handler callbacks with queue destroy. dynamicScratchState |= ERROR_HANDLER_TERMINATE; - HSA::hsa_signal_store_screlease(amd_queue_.queue_inactive_signal, 0x8000000000000000ull); while ((dynamicScratchState & ERROR_HANDLER_DONE) != ERROR_HANDLER_DONE) { + HSA::hsa_signal_store_screlease(amd_queue_.queue_inactive_signal, 0x8000000000000000ull); HSA::hsa_signal_wait_relaxed(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE, 0x8000000000000000ull, -1ull, HSA_WAIT_STATE_BLOCKED); - HSA::hsa_signal_store_relaxed(amd_queue_.queue_inactive_signal, 0x8000000000000000ull); } // Remove kfd exception handler @@ -1028,6 +1027,9 @@ bool AqlQueue::ExceptionHandler(hsa_signal_value_t error_code, void* arg) { if (queue->errors_callback_ != nullptr) { queue->errors_callback_(errorCode, queue->public_handle(), queue->errors_data_); } + Signal* signal = queue->exception_signal_; + queue->exceptionState = ERROR_HANDLER_DONE; + signal->StoreRelease(0); return false; }