From 7890b0559e574851179ed260ec3e6d4e0ea4107e Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Mon, 28 Nov 2022 14:03:51 +0530
Subject: [PATCH] SWDEV-348684 - Fixed catch-2 test failure (#3073)

Change-Id: Ib21c26a3c8e51f0b7f132ecd4ad209dfef1d0689
---
 .../streamperthread/hipStreamPerThrdTsts.cc   | 20 ++++++++++++-------
 1 file changed, 13 insertions(+), 7 deletions(-)

diff --git a/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/catch/unit/streamperthread/hipStreamPerThrdTsts.cc
index 95f5c27fd5..dec35e100c 100644
--- a/catch/unit/streamperthread/hipStreamPerThrdTsts.cc
+++ b/catch/unit/streamperthread/hipStreamPerThrdTsts.cc
@@ -47,9 +47,7 @@ THE SOFTWARE.
 #endif
 
 #include <hip_test_common.hh>
-#ifdef HT_AMD
-  #include "hip/hip_cooperative_groups.h"
-#endif
+#include "hip/hip_cooperative_groups.h"
 using namespace std::chrono;
 using namespace cooperative_groups;
 #if HT_AMD
@@ -69,8 +67,12 @@ __global__ void StreamPerThrd(int *Ad, int *Ad1, size_t n, int Pk_Clk,
     int64_t GpuFrq = (Pk_Clk * 1000);
     int64_t StrtTck = clock64();
     if (index == 0) {
-      // The following while loop checks the value in ptr for around 4 seconds
-      while ((clock64() - StrtTck) <= (6 * GpuFrq)) {
+      // The following while loop holds the execution for ~2 seconds.
+      // Busy sleep on nvidia
+      while ((clock64() - StrtTck) <= (2 * GpuFrq)) {
+        #if HT_AMD
+          __builtin_amdgcn_s_sleep(10);
+        #endif
       }
       if (WaitEvnt == 1) {
         *Ad1 = 1;
@@ -83,8 +85,12 @@ __global__ void StreamPerThrd(int *Ad, int *Ad1, size_t n, int Pk_Clk,
 __global__ void StreamPerThrd1(int *A, int Pk_Clk) {
   int64_t GpuFrq = (Pk_Clk * 1000);
   int64_t StrtTck = clock64();
-  // The following while loop checks the value in ptr for around 3-4 seconds
-  while ((clock64() - StrtTck) <= (3 * GpuFrq)) {
+  // The following while loop holds the execution for ~1 second
+  // Busy sleep on nvidia
+  while ((clock64() - StrtTck) <= (GpuFrq)) {
+    #if HT_AMD
+      __builtin_amdgcn_s_sleep(10);
+    #endif
   }
   *A = 1;
 }