From 3d6a18b67c1579a5c2e6e52d176fb0fd012a149b 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 --- runtime/hsa-runtime/core/runtime/runtime.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/runtime/hsa-runtime/core/runtime/runtime.cpp b/runtime/hsa-runtime/core/runtime/runtime.cpp index d6fae84346..51dcb75694 100644 --- a/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/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]);