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: 700f1d9abd]

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;
 }

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;