From bf79c21ea833334d56352e18997a0b25d65dd172 Mon Sep 17 00:00:00 2001
From: Brandon Potter <brandon.potter@amd.com>
Date: Tue, 10 Sep 2024 09:34:45 -0700
Subject: [PATCH] Conservatively use SEQ_CST atomics in IPC conduit

[ROCm/rocshmem commit: 7411c4559151e963c173b89910b89ffb3eb03b35]
---
 projects/rocshmem/src/ipc_policy.hpp | 22 +++++++++++-----------
 1 file changed, 11 insertions(+), 11 deletions(-)

diff --git a/projects/rocshmem/src/ipc_policy.hpp b/projects/rocshmem/src/ipc_policy.hpp
index c0190198ca..48b7b7fd40 100644
--- a/projects/rocshmem/src/ipc_policy.hpp
+++ b/projects/rocshmem/src/ipc_policy.hpp
@@ -68,34 +68,34 @@ class IpcOnImpl {
 
   template <typename T>
   __device__ T ipcAMOFetchAdd(T *val, T value) {
-    return __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED,
-                                  __HIP_MEMORY_SCOPE_AGENT);
+    return __hip_atomic_fetch_add(val, value, __ATOMIC_SEQ_CST,
+                                  __HIP_MEMORY_SCOPE_SYSTEM);
   }
 
   template <typename T>
   __device__ T ipcAMOFetchCas(T *val, T cond, T value) {
-    __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED,
-                                         __ATOMIC_RELAXED,
-                                         __HIP_MEMORY_SCOPE_AGENT);
+    __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_SEQ_CST,
+                                         __ATOMIC_SEQ_CST,
+                                         __HIP_MEMORY_SCOPE_SYSTEM);
     return cond;
   }
 
   template <typename T>
   __device__ void ipcAMOAdd(T *val, T value) {
-    __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED,
-                           __HIP_MEMORY_SCOPE_AGENT);
+    __hip_atomic_fetch_add(val, value, __ATOMIC_SEQ_CST,
+                           __HIP_MEMORY_SCOPE_SYSTEM);
   }
 
   template <typename T>
   __device__ void ipcAMOCas(T *val, T cond, T value) {
-    __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED,
-                                         __ATOMIC_RELAXED,
-                                         __HIP_MEMORY_SCOPE_AGENT);
+    __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_SEQ_CST,
+                                         __ATOMIC_SEQ_CST,
+                                         __HIP_MEMORY_SCOPE_SYSTEM);
   }
 
   template <typename T>
   __device__ void ipcAMOSet(T *val, T value) {
-    __hip_atomic_store(val, value, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+    __hip_atomic_store(val, value, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SYSTEM);
   }
 
   __device__ void zero_byte_read(int pe) {