SWDEV-419288 - Improve TstCoherency()

Fix originally wrong logic.

Change-Id: Ia4d3d9ee3d5712ce2e8914afcc2846b88306c483

catch/unit/memory/hipMemCoherencyTst.cc

@@ -35,8 +35,9 @@
 #include <chrono>
 
 __global__ void CoherentTst(int* ptr) {  // ptr was set to 1
-  atomicAdd(ptr, 1);                     // now ptr is 2
-  while (atomicCAS(ptr, 3, 4) != 3) {    // wait till ptr is 3, then change it to 4
+  atomicAdd_system(ptr, 1);              // now ptr is 2
+  while (atomicCAS_system(ptr, 3, 4) != 3) {
+    // wait till ptr is updated to 3 in host, then change it to 4
   }
 }
 
@@ -45,8 +46,6 @@ __global__  void SquareKrnl(int *ptr) {
   *ptr = (*ptr) * (*ptr);
 }
 
-
-
 // The variable below will work as signal to decide pass/fail
 static bool YES_COHERENT = false;
 
@@ -66,18 +65,28 @@ static void TstCoherency(int* ptr, bool hmmMem) {
   } else {
     CoherentTst<<<1, 1, 0, stream>>>(ptr);
   }
-
-  std::chrono::steady_clock::time_point start = std::chrono::steady_clock::now();
-  while (std::chrono::duration_cast<std::chrono::seconds>(std::chrono::steady_clock::now() - start)
-                 .count() < 3 &&
-         *ptr == 2) {
-  }          // wait till ptr is 2 from kernel or 3 seconds
-
-  *ptr += 1; // increment it to 3
+  // To prevent Windows batch dispatching issue, run inspecting code in thread
+  std::thread my_thread([ptr] {
+    int d = 0;
+    std::chrono::steady_clock::time_point start = std::chrono::steady_clock::now();
+    while (
+        std::chrono::duration_cast<std::chrono::seconds>(std::chrono::steady_clock::now() - start)
+            .count() <= 3) {
+      d = __sync_fetch_and_add(ptr, 0);  // Retrieve *ptr
+      if (d == 2) break; // If kernel has updated *ptr to 2, exit
+    }  // wait till ptr is updated to 2 from kernel or 3 seconds
+    if (d != 2) {
+      // 3 seconds should be long enough for kernel to update ptr
+      fprintf(stderr, "d = %d hasn't been updated to 2 in 3s\n", d);
+      return;
+    }
+    // increment it to 3
+    __sync_fetch_and_add(ptr, 1);
+  });
 
   HIP_CHECK(hipStreamSynchronize(stream));
   HIP_CHECK(hipStreamDestroy(stream));
-
+  my_thread.join();
   if (*ptr == 4) {
     YES_COHERENT = true;
   }