From 8d4608ed0ebf686cf035e79d41c240980c22d6ad Mon Sep 17 00:00:00 2001
From: Jay Cornwall <jay.cornwall@amd.com>
Date: Wed, 2 Jun 2021 16:59:09 -0500
Subject: [PATCH] Clear queue error code when not handling exceptions

If left non-zero the event loop will keep reinvoking the callback,
preventing AqlQueue::ExceptionHandler from running.

Change-Id: If85fbaf62f04ffd327ecf9d649aa23afad4442ce
---
 runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp | 3 +++
 1 file changed, 3 insertions(+)

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) {