From 28ef3b4b9eef177624ffc6a853f47f6a406fef32 Mon Sep 17 00:00:00 2001
From: Julia Jiang <56359287+jujiang-del@users.noreply.github.com>
Date: Fri, 23 Jan 2026 11:26:27 -0500
Subject: [PATCH] swdev-564412 - fix seg faults in atomic operation tests with
 muti-gpus (#2541)

---
 projects/hip-tests/catch/unit/atomics/bitwise_common.hh | 6 ++++--
 projects/hip-tests/catch/unit/atomics/min_max_common.hh | 2 ++
 2 files changed, 6 insertions(+), 2 deletions(-)

diff --git a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
index 23e6d9cf78..75ae13ab09 100644
--- a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
@@ -263,11 +263,12 @@ void TestCore(const TestParams& p) {
   // Test Values on Device
   TestType test_value = GetTestValue<TestType, operation>();
   for (auto i = 0u; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
     TestType* const mem_ptr =
         p.alloc_type == LinearAllocs::hipMalloc ? mem_devs[i].ptr() : mem_devs[i].host_ptr();
     HIP_CHECK(hipMemset(mem_ptr, 0, mem_alloc_size));
-    for (int i = 0; i < p.width * p.pitch / sizeof(TestType); ++i) {
-      HIP_CHECK(hipMemcpy(&mem_ptr[i], &test_value, sizeof(TestType), hipMemcpyHostToDevice));
+    for (int j = 0; j < p.width * p.pitch / sizeof(TestType); ++j) {
+      HIP_CHECK(hipMemcpy(&mem_ptr[j], &test_value, sizeof(TestType), hipMemcpyHostToDevice));
     }
   }
   // Launch Kernel and get back old vals
@@ -286,6 +287,7 @@ void TestCore(const TestParams& p) {
   }
   // Copy results back to Host
   for (auto i = 0u; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
     const auto device_offset = i * p.kernel_count * p.ThreadCount();
     HIP_CHECK(hipMemcpy(old_vals.data() + device_offset, old_vals_devs[i].ptr(),
                         old_vals_alloc_size, hipMemcpyDeviceToHost));
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 3df6658824..85d223f588 100644
--- a/projects/hip-tests/catch/unit/atomics/min_max_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/min_max_common.hh
@@ -286,6 +286,7 @@ void TestCore(const TestParams& p) {
   TestType test_value =
       std::is_floating_point_v<TestType> ? kFloatingPointTestValue : kIntegerTestValue;
   for (auto i = 0u; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
     TestType* const mem_ptr =
       p.alloc_type == LinearAllocs::hipMalloc ? mem_devs[i].ptr() : mem_devs[i].host_ptr();
     HIP_CHECK(hipMemset(mem_ptr, 0, mem_alloc_size));
@@ -310,6 +311,7 @@ void TestCore(const TestParams& p) {
   }
   // Copy Results back to Host
   for (auto i = 0u; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
     const auto device_offset = i * p.kernel_count * p.ThreadCount();
     HIP_CHECK(hipMemcpy(old_vals.data() + device_offset, old_vals_devs[i].ptr(),
                         old_vals_alloc_size, hipMemcpyDeviceToHost));