From 6dc8a4ae556a6728f42683d464c18dd58959e31e Mon Sep 17 00:00:00 2001
From: Jaydeep Patel <jaydeepkumar.patel@amd.com>
Date: Tue, 5 Nov 2024 17:56:24 +0000
Subject: [PATCH] rocr: Decrement counter only if event is popped

Also restore dead signals cleanup for old path when HSA_WAIT_ANY_DEBUG
is used.

Change-Id: I51a7404991443c9f6cbf57b4b9e9faa694b9538c


[ROCm/ROCR-Runtime commit: 700f1d9abd44598047c7352a0cee40f2e16bf4f1]
---
 .../runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp       | 3 +++
 .../runtime/hsa-runtime/core/runtime/runtime.cpp             | 5 +++--
 2 files changed, 6 insertions(+), 2 deletions(-)

diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
index ea6b34d43b..8abec94838 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
@@ -1014,6 +1014,9 @@ static bool GangCopyCompleteHandler(hsa_signal_value_t, void *arg ) {
   core::Signal *gang_signal = reinterpret_cast<core::Signal*>(arg);
   if (gang_signal->IsValid()) {
     gang_signal->DestroySignal();
+    if (!gang_signal->IsValid()) {
+      return false;
+    }
   }
   return true;
 }
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 4fea61c544..106386b13b 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp
@@ -1684,8 +1684,9 @@ void Runtime::AsyncEventsLoop(void* _eventsInfo) {
             if (i == 0) {
               hsa_signal_handle(async_events_control_.wake)->StoreRelaxed(0);
             } else {
-              processEvent(i, value);
-              i--;
+              if (!processEvent(i, value)) {
+                i--;
+              }
             }
             if (!wait_any) {
               finish = true;