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