From 6b398eb72cc6c70fb776bec93efebb7578bcc981 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Thu, 27 May 2021 15:41:31 -0500 Subject: [PATCH] Improve async handler performance. Under high async handler load signal retention and event sorting become bottlenecks. This change processes more handlers in a single pass to amortize wait_any overheads. Change-Id: I8b276e102db647e3858e120547aa0c6fca85ab4c --- runtime/hsa-runtime/core/runtime/runtime.cpp | 44 ++++++++++++++++---- 1 file changed, 36 insertions(+), 8 deletions(-) diff --git a/runtime/hsa-runtime/core/runtime/runtime.cpp b/runtime/hsa-runtime/core/runtime/runtime.cpp index a6bb6cb853..d6fae84346 100644 --- a/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -1051,14 +1051,42 @@ void Runtime::AsyncEventsLoop(void*) { if (index == 0) { hsa_signal_handle(async_events_control_.wake)->StoreRelaxed(0); } else if (index != -1) { - // No error or timout occured, process the handler - 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(); + // No error or timout occured, process the handlers + for (size_t i = index; i < async_events_.Size(); i++) { + hsa_signal_handle sig(async_events_.signal_[i]); + + value = atomic::Load(&sig->signal_.value, std::memory_order_relaxed); + bool condition_met = false; + + switch (async_events_.cond_[i]) { + case HSA_SIGNAL_CONDITION_EQ: { + condition_met = (value == async_events_.value_[i]); + break; + } + case HSA_SIGNAL_CONDITION_NE: { + condition_met = (value != async_events_.value_[i]); + break; + } + case HSA_SIGNAL_CONDITION_GTE: { + condition_met = (value >= async_events_.value_[i]); + break; + } + case HSA_SIGNAL_CONDITION_LT: { + condition_met = (value < async_events_.value_[i]); + break; + } + } + + if (condition_met) { + assert(async_events_.handler_[i] != NULL); + bool keep = async_events_.handler_[i](value, async_events_.arg_[i]); + if (!keep) { + hsa_signal_handle(async_events_.signal_[i])->Release(); + async_events_.CopyIndex(i, async_events_.Size() - 1); + async_events_.PopBack(); + i--; + } + } } }