diff --git a/catch/unit/memory/hipMemCoherencyTst.cc b/catch/unit/memory/hipMemCoherencyTst.cc index a5dbaed656..6e2c2b588d 100644 --- a/catch/unit/memory/hipMemCoherencyTst.cc +++ b/catch/unit/memory/hipMemCoherencyTst.cc @@ -35,8 +35,9 @@ #include __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::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::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; }