From 938994cd60c7da67d23f938a16efbd6581a073ea Mon Sep 17 00:00:00 2001
From: Jatin Chaudhary <JatinJaikishan.Chaudhary@amd.com>
Date: Tue, 12 Dec 2023 15:25:54 +0000
Subject: [PATCH] SWDEV-431399 - fix randomly failing timing based tests

use blocking context instead of time based kernel launch

Change-Id: Ib4cdf913f18c80ce8173317811c0e6e3eb5acb76


[ROCm/hip-tests commit: c3646564206c89c904086ac8647aab844b8b356f]
---
 .../hip-tests/catch/unit/memory/hipFree.cc    | 29 ++++++++++++++-----
 1 file changed, 22 insertions(+), 7 deletions(-)

diff --git a/projects/hip-tests/catch/unit/memory/hipFree.cc b/projects/hip-tests/catch/unit/memory/hipFree.cc
index a9e0876f9d..23b45d5139 100644
--- a/projects/hip-tests/catch/unit/memory/hipFree.cc
+++ b/projects/hip-tests/catch/unit/memory/hipFree.cc
@@ -46,7 +46,6 @@ enum class FreeType { DevFree, ArrayFree, ArrayDestroy, HostFree };
 
 // Amount of time kernel should wait
 using namespace std::chrono_literals;
-const std::chrono::duration<uint64_t, std::milli> delay = 50ms;
 constexpr size_t numAllocs = 10;
 
 TEST_CASE("Unit_hipFreeImplicitSyncDev") {
@@ -54,9 +53,14 @@ TEST_CASE("Unit_hipFreeImplicitSyncDev") {
   size_t size_mult = GENERATE(1, 32, 64, 128, 256);
   HIP_CHECK(hipMalloc(&devPtr, sizeof(*devPtr) * size_mult));
 
-  LaunchDelayKernel(delay);
-  // make sure device is busy
+  HipTest::BlockingContext b_context{nullptr};
+
+  b_context.block_stream();
+  REQUIRE(b_context.is_blocked());
+
   HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady);
+  b_context.unblock_stream();
+
   HIP_CHECK(hipFree(devPtr));
   HIP_CHECK(hipStreamQuery(nullptr));
 }
@@ -67,9 +71,14 @@ TEST_CASE("Unit_hipFreeImplicitSyncHost") {
 
   HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(*hostPtr) * size_mult));
 
-  LaunchDelayKernel(delay);
-  // make sure device is busy
+  HipTest::BlockingContext b_context{nullptr};
+
+  b_context.block_stream();
+  REQUIRE(b_context.is_blocked());
+
   HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady);
+  b_context.unblock_stream();
+
   HIP_CHECK(hipHostFree(hostPtr));
   HIP_CHECK(hipStreamQuery(nullptr));
 }
@@ -79,6 +88,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo
   CHECK_IMAGE_SUPPORT
 
   using vec_info = vector_info<TestType>;
+  const std::chrono::duration<uint64_t, std::milli> delay = 50ms;
   DriverContext ctx;
 
 
@@ -124,9 +134,14 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo
   hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
 
   HIP_CHECK(hipMallocArray(&arrayPtr, &desc, extent.width, extent.height, hipArrayDefault));
-  LaunchDelayKernel(delay);
-  // make sure device is busy
+  HipTest::BlockingContext b_context{nullptr};
+
+  b_context.block_stream();
+  REQUIRE(b_context.is_blocked());
+
   HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady);
+  b_context.unblock_stream();
+
   // Second free segfaults
   SECTION("ArrayDestroy") {
     HIP_CHECK(hipArrayDestroy(arrayPtr));