SWDEV-516488 - Sync all devices. (#349)

[ROCm/hip-tests commit: 863f8828c3]

projects/hip-tests/catch/unit/atomics/arithmetic_common.hh

@@ -443,7 +443,10 @@ 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());
+  for (auto i = 0; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
+    HIP_CHECK(hipDeviceSynchronize());
+  }
   // Copy results back to Host
   for (auto i = 0u; i < p.num_devices; ++i) {
     HIP_CHECK(hipSetDevice(i));

projects/hip-tests/catch/unit/atomics/atomicExch_common.hh

@@ -229,7 +229,10 @@ 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 = 0; i < p.num_devices; ++i) {
+      HIP_CHECK(hipSetDevice(i));
+      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(),

projects/hip-tests/catch/unit/atomics/bitwise_common.hh

@@ -279,7 +279,10 @@ void TestCore(const TestParams& p) {
                                                                       old_vals);
     }
   }
-  HIP_CHECK(hipDeviceSynchronize());
+  for (auto i = 0; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
+    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();

projects/hip-tests/catch/unit/atomics/min_max_common.hh

@@ -309,8 +309,10 @@ void TestCore(const TestParams& p) {
                                                                       old_vals);
     }
   }
-  HIP_CHECK(hipDeviceSynchronize());
-
+  for (auto i = 0; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
+    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();