SWDEW_516488 - Fix possible atomics tests' failures (#292)

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));

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);
     }
   }
-}
+}

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();

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) {