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 __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 __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 __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 __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 __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) {