From 7b760bb023b617f4d1becf8a0d96f3044cadb69c Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 22 Oct 2024 17:22:51 +0000 Subject: [PATCH] fix barrier synchronization on gfx90a [ROCm/rocshmem commit: 87db7f7d387b004abf171671a51f1aed9b28c91e] --- projects/rocshmem/src/ipc/context_ipc_device.cpp | 1 + projects/rocshmem/src/ipc/context_ipc_device_coll.cpp | 3 +++ 2 files changed, 4 insertions(+) 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