From cf65854e7986378f4735ff31b9eb0197a53a4192 Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Sun, 9 Jul 2023 09:08:35 +0530
Subject: [PATCH] SWDEV-407072 - Fix sync issues seen on some event tests
 (#358)

Change-Id: Ib904b9cca7329abcd02004daaeafd170014a1ec0
---
 catch/unit/event/Unit_hipEventElapsedTime.cc |  7 +-
 catch/unit/event/Unit_hipEventQuery.cc       | 70 ++------------------
 2 files changed, 12 insertions(+), 65 deletions(-)

diff --git a/catch/unit/event/Unit_hipEventElapsedTime.cc b/catch/unit/event/Unit_hipEventElapsedTime.cc
index 7149c8cbea..93deb7f842 100644
--- a/catch/unit/event/Unit_hipEventElapsedTime.cc
+++ b/catch/unit/event/Unit_hipEventElapsedTime.cc
@@ -158,7 +158,9 @@ TEST_CASE("Unit_hipEventElapsedTime_NotReady_Negative") {
   // Record start event
   HIP_CHECK(hipEventRecord(start, nullptr));
 
-  HipTest::runKernelForDuration(std::chrono::milliseconds(1000));
+  HipTest::BlockingContext b_context{nullptr};
+  b_context.block_stream();  // blocked stream
+  REQUIRE(b_context.is_blocked());
 
   // Record stop event
   HIP_CHECK(hipEventRecord(stop, nullptr));
@@ -166,7 +168,8 @@ TEST_CASE("Unit_hipEventElapsedTime_NotReady_Negative") {
   // stop event has not been completed
   float tElapsed = 1.0f;
   HIP_CHECK_ERROR(hipEventQuery(stop), hipErrorNotReady);
-  HIP_ASSERT(hipEventElapsedTime(&tElapsed,start,stop) == hipErrorNotReady);
+  HIP_ASSERT(hipEventElapsedTime(&tElapsed, start, stop) == hipErrorNotReady);
+  b_context.unblock_stream();
 
   HIP_CHECK(hipStreamSynchronize(nullptr));
   HIP_CHECK(hipEventDestroy(start));
diff --git a/catch/unit/event/Unit_hipEventQuery.cc b/catch/unit/event/Unit_hipEventQuery.cc
index 726a57ca53..b9606fccf4 100644
--- a/catch/unit/event/Unit_hipEventQuery.cc
+++ b/catch/unit/event/Unit_hipEventQuery.cc
@@ -19,57 +19,6 @@ THE SOFTWARE.
 
 #include <hip_test_common.hh>
 
-/**
- * @addtogroup hipEventQuery hipEventQuery
- * @{
- * @ingroup EventTest
- * `hipEventQuery(hipEvent_t event)` -
- * Query the status of the specified event.
- * ________________________
- * Test cases from other modules:
- *  - @ref Unit_hipEventIpc
- */
-
-// Since we can not use atomic*_system on every gpu, we will use wait based on clock rate.
-// This wont be accurate since current clock rate of a GPU varies depending on many variables
-// including thermals, load, utilization etc
-__global__ void waitKernel(int clockRate, int seconds) {
-  auto start = clock();
-  auto ms = seconds * 1000;
-  long long waitTill = clockRate * (long long)ms;
-  while (1) {
-    auto end = clock();
-    if ((end - start) > waitTill) {
-      return;
-    }
-  }
-}
-
-__global__ void waitKernel_gfx11(int clockRate, int seconds) {
-#if HT_AMD
-  auto start = wall_clock64();
-  auto ms = seconds * 1000;
-  long long waitTill = clockRate * (long long)ms;
-  while (1) {
-    auto end = wall_clock64();
-    if ((end - start) > waitTill) {
-      return;
-    }
-  }
-#endif
-}
-
-/**
- * Test Description
- * ------------------------
- *  - Query events with a single and with multiple devices.
- * Test source
- * ------------------------
- *  - unit/event/Unit_hipEventQuery.cc
- * Test requirements
- * ------------------------
- *  - HIP_VERSION >= 5.2
- */
 TEST_CASE("Unit_hipEventQuery_DifferentDevice") {
   hipEvent_t event1{}, event2{};
   HIP_CHECK(hipEventCreate(&event1));
@@ -81,19 +30,14 @@ TEST_CASE("Unit_hipEventQuery_DifferentDevice") {
   HIP_CHECK(hipStreamCreate(&stream));
   REQUIRE(stream != nullptr);
 
-  hipDeviceProp_t prop{};
-  HIP_CHECK(hipGetDeviceProperties(&prop, 0));
-  auto clockRate = prop.clockRate;
-
-  // Start kernel on 1st device
+  HipTest::BlockingContext b_context1{stream};  // og context
+  // Block stream
   {
     HIP_CHECK(hipSetDevice(0));
     HIP_CHECK(hipEventRecord(event1, stream));
 
-    auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel;
-    // Start kernel and wait for 3 seconds
-    // Make sure you increase this time if you add more tests here
-    waitKernel_used<<<1, 1, 0, stream>>>(clockRate, 3);
+    b_context1.block_stream();  // blocked stream
+    REQUIRE(b_context1.is_blocked());
 
     HIP_CHECK(hipEventRecord(event2, stream));
 
@@ -101,7 +45,7 @@ TEST_CASE("Unit_hipEventQuery_DifferentDevice") {
     HIP_CHECK(hipEventQuery(event1));  // Should be done
 
     HIP_CHECK_ERROR(hipEventQuery(event2),
-                    hipErrorNotReady);  // Wont be done since kernel is waiting
+                    hipErrorNotReady);  // Wont be done since stream is blocked
   }
 
   // If other devices are available, set it
@@ -114,8 +58,8 @@ TEST_CASE("Unit_hipEventQuery_DifferentDevice") {
     HIP_CHECK(hipEventQuery(event1));
     HIP_CHECK_ERROR(hipEventQuery(event2), hipErrorNotReady);
 
-    // Sync
-    // if it hangs here it means GPU kernel cant see the value update by atomic store
+    b_context1.unblock_stream();
+
     HIP_CHECK(hipEventSynchronize(event2));
 
     // Query, should be done now