diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 3a9ab56cb3..6bb4788a05 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -921,6 +921,9 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { errorCode = HSA_STATUS_ERROR; fatal = true; } + } else { + // Not handling exceptions, clear so that ExceptionHandler can run. + HSA::hsa_signal_store_relaxed(queue->amd_queue_.queue_inactive_signal, 0); } if (errorCode == HSA_STATUS_SUCCESS) {