From 91681bfd287ceb8005805bce658a0007ff78442c Mon Sep 17 00:00:00 2001 From: "Sang, Tao" Date: Wed, 2 Jul 2025 00:35:40 -0400 Subject: [PATCH] SWDEV-516488 - Sync all devices. (#349) [ROCm/hip-tests commit: 863f8828c3a99198842d455044522a2a41e24d0c] --- projects/hip-tests/catch/unit/atomics/arithmetic_common.hh | 5 ++++- projects/hip-tests/catch/unit/atomics/atomicExch_common.hh | 5 ++++- projects/hip-tests/catch/unit/atomics/bitwise_common.hh | 5 ++++- projects/hip-tests/catch/unit/atomics/min_max_common.hh | 6 ++++-- 4 files changed, 16 insertions(+), 5 deletions(-) diff --git a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh index b0725a4e71..add775a279 100644 --- a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh +++ b/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(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)); diff --git a/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh b/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh index 92db58c210..e0fcf84c87 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh +++ b/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(), diff --git a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh index e646c05a58..b611679683 100644 --- a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh +++ b/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(); 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 09cd06e34c..467cfdcae1 100644 --- a/projects/hip-tests/catch/unit/atomics/min_max_common.hh +++ b/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();