fix barrier synchronization on gfx90a

[ROCm/rocshmem commit: 87db7f7d38]

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

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