From 976d0574cf48af14f133ebad359818f945ecef36 Mon Sep 17 00:00:00 2001
From: "Sang, Tao" <Tao.Sang@amd.com>
Date: Fri, 13 Jun 2025 20:03:25 -0400
Subject: [PATCH] SWDEW_516488 - Fix possible atomics tests' failures (#292)

[ROCm/hip-tests commit: 6c60fcafd870544e9108a0614b18ab037e45e75a]
---
 projects/hip-tests/catch/unit/atomics/arithmetic_common.hh | 2 +-
 projects/hip-tests/catch/unit/atomics/atomicExch_common.hh | 4 ++--
 projects/hip-tests/catch/unit/atomics/bitwise_common.hh    | 1 +
 projects/hip-tests/catch/unit/atomics/min_max_common.hh    | 1 +
 4 files changed, 5 insertions(+), 3 deletions(-)

diff --git a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
index 4fb2111d74..b0725a4e71 100644
--- a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
@@ -443,7 +443,7 @@ void TestCore(const TestParams& p) {
   mem_devs.emplace_back(LinearAllocs::hipHostMalloc, mem_alloc_size);
   PerformHostAtomicOperation<TestType, operation>(p, mem_devs[p.num_devices].host_ptr(), old_vals.data());
 
-
+  HIP_CHECK(hipDeviceSynchronize());
   // Copy results back to Host
   for (auto i = 0u; i < p.num_devices; ++i) {
     HIP_CHECK(hipSetDevice(i));
diff --git a/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh b/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh
index fcb1fc770e..92db58c210 100644
--- a/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh
@@ -229,7 +229,7 @@ class AtomicExchCRTP {
 
     PerformHostAtomicExchange(p.host_thread_count, host_iters_per_thread, mem_dev.host_ptr(),
                               old_vals.data(), p);
-
+    HIP_CHECK(hipDeviceSynchronize());
     for (auto i = 0u; i < p.num_devices; ++i) {
       const auto device_offset = i * p.kernel_count * thread_count;
       HIP_CHECK(hipMemcpy(old_vals.data() + device_offset, old_vals_devs[i].ptr(),
@@ -419,4 +419,4 @@ void AtomicExchMultipleDeviceMultipleKernelAndHostTest(const unsigned int num_de
       AtomicExch<TestType, false, AtomicScopes::system>().run(params);
     }
   }
-}
\ No newline at end of file
+}
diff --git a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
index cabbdaf13a..e646c05a58 100644
--- a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
@@ -279,6 +279,7 @@ void TestCore(const TestParams& p) {
                                                                       old_vals);
     }
   }
+  HIP_CHECK(hipDeviceSynchronize());
   // Copy results back to Host
   for (auto i = 0u; i < p.num_devices; ++i) {
     const auto device_offset = i * p.kernel_count * p.ThreadCount();
diff --git a/projects/hip-tests/catch/unit/atomics/min_max_common.hh b/projects/hip-tests/catch/unit/atomics/min_max_common.hh
index 8f7c21cb07..09cd06e34c 100644
--- a/projects/hip-tests/catch/unit/atomics/min_max_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/min_max_common.hh
@@ -309,6 +309,7 @@ void TestCore(const TestParams& p) {
                                                                       old_vals);
     }
   }
+  HIP_CHECK(hipDeviceSynchronize());
 
   // Copy Results back to Host
   for (auto i = 0u; i < p.num_devices; ++i) {