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: 3d6a18b67c]

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]);