fixed barrier issue on MI250X

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();