From 437ce0b8df73d922c545b5303d34c8cc1fdc265b Mon Sep 17 00:00:00 2001
From: Victor Zhang <111778801+victzhan@users.noreply.github.com>
Date: Sun, 2 Nov 2025 21:45:44 -0500
Subject: [PATCH] fix atomics SystemTest() use after free (#1595)

---
 .../hip-tests/catch/unit/atomics/memory_order_common.hh  | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/projects/hip-tests/catch/unit/atomics/memory_order_common.hh b/projects/hip-tests/catch/unit/atomics/memory_order_common.hh
index 8251f8e790..5ecabd1925 100644
--- a/projects/hip-tests/catch/unit/atomics/memory_order_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/memory_order_common.hh
@@ -259,6 +259,9 @@ template <BuiltinAtomicOperation operation, int memory_order> void SystemTest()
       });
       ConsumerKernel<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>
           <<<1, 1>>>(flag.ptr(), data.ptr(), ret.ptr());
+
+      HIP_CHECK(hipDeviceSynchronize());
+      host_thread.join();
     }
 
     SECTION("Device producer - Host consumer") {
@@ -268,12 +271,12 @@ template <BuiltinAtomicOperation operation, int memory_order> void SystemTest()
       });
       ProducerKernel<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>
           <<<1, 1>>>(flag.ptr(), data.ptr());
+
+      HIP_CHECK(hipDeviceSynchronize());
+      host_thread.join();
     }
   }
 
-  HIP_CHECK(hipDeviceSynchronize());
-  host_thread.join();
-
   REQUIRE(ret.ptr()[0] == kTestValue);
 }