AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source
This commit is contained in:
Sang, Tao
2025-07-02 00:35:40 -04:00
committed by GitHub
parent 9499bf6d1e
commit 863f8828c3
4 changed files with 16 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/atomics/arithmetic_common.hh
+4 -1
View File
@@ -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));
catch/unit/atomics/atomicExch_common.hh
+4 -1
View File
@@ -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(),
catch/unit/atomics/bitwise_common.hh
+4 -1
View File
@@ -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();
catch/unit/atomics/min_max_common.hh
+4 -2
View File
@@ -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();