swdev-564412 - fix seg faults in atomic operation tests with muti-gpus (#2541)

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

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