From e2f6a65284bc335ef9a684e21bbf02828b17433e Mon Sep 17 00:00:00 2001
From: Yiltan Hassan Temucin <yiltan.temucin@amd.com>
Date: Wed, 2 Oct 2024 09:20:07 -0700
Subject: [PATCH] fixed barrier issue on MI250X

---
 src/ipc/context_ipc_device_coll.cpp | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/src/ipc/context_ipc_device_coll.cpp b/src/ipc/context_ipc_device_coll.cpp
index 804be2efbf..1465b8ffe0 100644
--- a/src/ipc/context_ipc_device_coll.cpp
+++ b/src/ipc/context_ipc_device_coll.cpp
@@ -35,6 +35,9 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
   if (pe == PE_start) {
     // Go through all PE offsets (except current offset = 0)
     // and wait until they all reach
+#if defined(__gfx90a__)
+    __threadfence_system();
+#endif /* __gfx90a__ */
     for (size_t i = 1; i < n_pes; i++) {
       wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val);
       pSync[i] = ROC_SHMEM_SYNC_VALUE;
@@ -49,6 +52,9 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
     // Mark current PE offset as reached
     size_t pe_offset = (pe - PE_start) / stride;
     put_nbi(&pSync[pe_offset], &flag_val, 1, PE_start);
+#if defined(__gfx90a__)
+    __threadfence_system();
+#endif /* __gfx90a__ */
     wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
     pSync[0] = ROC_SHMEM_SYNC_VALUE;
     threadfence_system();