diff --git a/projects/rocshmem/src/ipc/context_ipc_device.cpp b/projects/rocshmem/src/ipc/context_ipc_device.cpp
index 275b15fe82..ee80694578 100644
--- a/projects/rocshmem/src/ipc/context_ipc_device.cpp
+++ b/projects/rocshmem/src/ipc/context_ipc_device.cpp
@@ -58,6 +58,7 @@ __host__ IPCContext::IPCContext(Backend *b)
 }
 
 __device__ void IPCContext::threadfence_system() {
+  __threadfence_system();
 }
 
 __device__ void IPCContext::ctx_create() {
diff --git a/projects/rocshmem/src/ipc/context_ipc_device_coll.cpp b/projects/rocshmem/src/ipc/context_ipc_device_coll.cpp
index 7d7f39919c..aa7e7be9ae 100644
--- a/projects/rocshmem/src/ipc/context_ipc_device_coll.cpp
+++ b/projects/rocshmem/src/ipc/context_ipc_device_coll.cpp
@@ -47,6 +47,9 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
     // Announce to other PEs that all have reached
     for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
       internal_putmem(&pSync[0], &flag_val, sizeof(*pSync), j);
+#if defined(__gfx90a__)
+        __threadfence_system();
+#endif /* __gfx90a__ */
     }
   } else {
     // Mark current PE offset as reached