From 183531963c3f86ee81c1694df13b3d365ffebae7 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Thu, 24 Jun 2021 02:32:13 -0500 Subject: [PATCH] Always execute the first satisfied async signal handler. Certain special signals do not carry their updates via their signal value. These signals are wrappers around special KFD events, of which the only current instance informs about VM faults. We either need to check each signal for this special event type or rely on the checking done in hsa_amd_signal_wait_any. Since there will always be a small number of these signals it doesn't make much since to penalize the performance path with this check. Additionally we know that the signal indicated by hsa_amd_signal_wait_any is satisfied so don't need to recheck it's conditions. Change-Id: I9fc6298300ad543d823ecd28ca8fab4ad26c23ef [ROCm/ROCR-Runtime commit: 3d6a18b67c1579a5c2e6e52d176fb0fd012a149b] --- .../runtime/hsa-runtime/core/runtime/runtime.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) 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 d6fae84346..51dcb75694 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -1052,6 +1052,15 @@ void Runtime::AsyncEventsLoop(void*) { hsa_signal_handle(async_events_control_.wake)->StoreRelaxed(0); } else if (index != -1) { // No error or timout occured, process the handlers + // Call handler for the known satisfied signal. + assert(async_events_.handler_[index] != NULL); + bool keep = async_events_.handler_[index](value, async_events_.arg_[index]); + if (!keep) { + hsa_signal_handle(async_events_.signal_[index])->Release(); + async_events_.CopyIndex(index, async_events_.Size() - 1); + async_events_.PopBack(); + } + // Check remaining signals before sleeping. for (size_t i = index; i < async_events_.Size(); i++) { hsa_signal_handle sig(async_events_.signal_[i]);