SWDEV-481514, SWDEV-482400 - Use correct way to add delay/wait on GPU.

Change-Id: I163896333bf741053173d636a499ed6406e77a09

catch/unit/module/hipExtModuleLaunchKernel.cc

@@ -346,7 +346,12 @@ void ModuleLaunchKernel::AllocateMemory() {
   HIP_CHECK(hipMemcpy(Ad, A, SIZE*sizeof(int), hipMemcpyHostToDevice));
   HIP_CHECK(hipMemcpy(Bd, B, SIZE*sizeof(int), hipMemcpyHostToDevice));
   int clkRate = 0;
+  #if HT_AMD
+  HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
+  #endif
+  #if HT_NVIDIA
   HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
+  #endif
   args1._Ad = Ad;
   args1._Bd = Bd;
   args1._Cd = C;

catch/unit/module/matmul.cpp

@@ -48,34 +48,32 @@ extern "C" __global__ void KernelandExtraParams(int* A, int* B, int* C,
   D[ROW * N + COL] = tmpSum;
 }
 
+__device__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
+  while (interval--) {
+    #if HT_AMD
+    uint64_t start = wall_clock64();
+    while (wall_clock64() - start < ticks_per_ms) {
+      __builtin_amdgcn_s_sleep(10);
+    }
+    #endif
+    #if HT_NVIDIA
+    uint64_t start = clock64();
+    while (clock64() - start < ticks_per_ms) {
+    }
+    #endif
+  }
+}
+
 extern "C" __global__ void SixteenSecKernel(int clockrate) {
-  uint64_t wait_t = 16000,
+  Delay(16000, clockrate);
-  start = clock64()/clockrate, cur;
-  do { cur = clock64()/clockrate-start;}while (cur < wait_t);
 }
 
 extern "C" __global__ void TwoSecKernel(int clockrate) {
-  if (deviceGlobal == 0x2222) {
+  Delay(2000, clockrate);
-    deviceGlobal = 0x3333;
-  }
-  uint64_t wait_t = 2000,
-  start = clock64()/clockrate, cur;
-  do { cur = clock64()/clockrate-start;}while (cur < wait_t);
-  if (deviceGlobal != 0x3333) {
-    deviceGlobal = 0x5555;
-  }
 }
 
 extern "C" __global__ void FourSecKernel(int clockrate) {
-  if (deviceGlobal == 1) {
+  Delay(4000, clockrate);
-    deviceGlobal = 0x2222;
-  }
-  uint64_t wait_t = 4000,
-  start = clock64()/clockrate, cur;
-  do { cur = clock64()/clockrate-start;}while (cur < wait_t);
-  if (deviceGlobal == 0x2222) {
-    deviceGlobal = 0x4444;
-  }
 }
 
 extern "C" __global__ void dummyKernel() {