From 5244a426c7c9f0360a322cae4ef6e47e080a6bc8 Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Mon, 9 Dec 2024 05:29:56 +0000 Subject: [PATCH] SWDEV-496544 - Memory needs visibility system wide as CPU and GPU both needs to see updated value by each other. Change-Id: Ia80e69455aeb06b9281f75313ba39b621d6929c5 --- catch/multiproc/hipMemCoherencyTstMProc.cc | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/catch/multiproc/hipMemCoherencyTstMProc.cc b/catch/multiproc/hipMemCoherencyTstMProc.cc index dec36e6dec..0de5b75cef 100644 --- a/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -44,10 +44,18 @@ __global__ void CoherentTst(int *ptr, int PeakClk) { // Incrementing the value by 1 int64_t GpuFrq = int64_t(PeakClk) * 1000; int64_t StrtTck = clock64(); + #if HT_AMD + atomicAdd_system(ptr, 1); + #else atomicAdd(ptr, 1); + #endif // The following while loop checks the value in ptr for around 3-4 seconds while ((clock64() - StrtTck) <= (3 * GpuFrq)) { + #if HT_AMD + if (atomicCAS_system(ptr, 3, 4) == 3) break; + #else if (atomicCAS(ptr, 3, 4) == 3) break; + #endif } } @@ -56,10 +64,10 @@ __global__ void CoherentTst_gfx11(int *ptr, int PeakClk) { // Incrementing the value by 1 int64_t GpuFrq = int64_t(PeakClk) * 1000; int64_t StrtTck = clock_function(); - atomicAdd(ptr, 1); + atomicAdd_system(ptr, 1); // The following while loop checks the value in ptr for around 3-4 seconds while ((clock_function() - StrtTck) <= (3 * GpuFrq)) { - if (atomicCAS(ptr, 3, 4) == 3) break; + if (atomicCAS_system(ptr, 3, 4) == 3) break; } #endif }