From dd8265fc877c33b2ba3ea9741bea8027b48845e5 Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Thu, 10 Jun 2021 20:30:32 -0400
Subject: [PATCH] SWDEV-290367 - Make sure HIP checks for GPU signal

Check GPU signal status for event before falling into CPU command
status validation

Change-Id: I66f15752d7dca550c0fa1a2252ec5a63817391c3
---
 hipamd/src/hip_event.cpp  | 20 +++++++++++++++-----
 hipamd/src/hip_stream.cpp |  7 ++++++-
 2 files changed, 21 insertions(+), 6 deletions(-)

diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp
index 09b73913cf..898b2406e7 100755
--- a/hipamd/src/hip_event.cpp
+++ b/hipamd/src/hip_event.cpp
@@ -35,8 +35,12 @@ bool Event::ready() {
   if (event_->status() != CL_COMPLETE) {
     event_->notifyCmdQueue();
   }
-
-  return (event_->status() == CL_COMPLETE);
+  // Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status
+  bool ready = g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_);
+  if (!ready) {
+    ready = (event_->status() == CL_COMPLETE);
+  }
+  return ready;
 }
 
 hipError_t Event::query() {
@@ -58,7 +62,11 @@ hipError_t Event::synchronize() {
     return hipSuccess;
   }
 
-  event_->awaitCompletion();
+  // Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status
+  static constexpr bool kWaitCompletion = true;
+  if (!g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, kWaitCompletion)) {
+    event_->awaitCompletion();
+  }
 
   return hipSuccess;
 }
@@ -84,8 +92,7 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) {
   }
   amd::ScopedLock stopLock(eStop.lock_);
 
-  if (event_ == nullptr ||
-      eStop.event_  == nullptr) {
+  if (event_ == nullptr || eStop.event_  == nullptr) {
     return hipErrorInvalidHandle;
   }
 
@@ -107,6 +114,9 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) {
     ms = static_cast<float>(static_cast<int64_t>(command->event().profilingInfo().end_) - time())/1000000.f;
     command->release();
   } else {
+    // Note: with direct dispatch eStop.ready() relies on HW event, but CPU status can be delayed.
+    // Hence for now make sure CPU status is updated by calling awaitCompletion();
+    eStop.event_->awaitCompletion();
     ms = static_cast<float>(eStop.time() - time())/1000000.f;
   }
   return hipSuccess;
diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp
index 801895cfd1..f35467e075 100755
--- a/hipamd/src/hip_stream.cpp
+++ b/hipamd/src/hip_stream.cpp
@@ -421,7 +421,12 @@ hipError_t hipStreamQuery(hipStream_t stream) {
   if (command->type() != 0) {
     event.notifyCmdQueue();
   }
-  hipError_t status = (command->status() == CL_COMPLETE) ? hipSuccess : hipErrorNotReady;
+  // Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status
+  bool ready = command->queue()->device().IsHwEventReady(event);
+  if (!ready) {
+    ready = (command->status() == CL_COMPLETE);
+  }
+  hipError_t status = ready ? hipSuccess : hipErrorNotReady;
   command->release();
   HIP_RETURN(status);
 }