From cb14fccdccacb7598192d886dc639ba94791ea70 Mon Sep 17 00:00:00 2001
From: Larry Meadows <lmeadows@amd.com>
Date: Thu, 25 Sep 2025 14:07:07 -0700
Subject: [PATCH]   - LL Protocol: Add missing fences for gfx950, this fixes
 the hang issue (#1932)

- Remove asm flat_store_dwordx4, not needed
---
 src/device/prims_ll.h | 18 +++++-------------
 1 file changed, 5 insertions(+), 13 deletions(-)

diff --git a/src/device/prims_ll.h b/src/device/prims_ll.h
index 0995730fd3..d5867ef4b1 100644
--- a/src/device/prims_ll.h
+++ b/src/device/prims_ll.h
@@ -261,15 +261,6 @@ private:
 
   __device__ void storeLL(union ncclLLFifoLine* dst, uint64_t val, uint32_t flag) {
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
-#if (defined(__gfx950__) && defined(HIP_HOST_UNCACHED_MEMORY))
-    using Vec = uint32_t __attribute__((ext_vector_type(4)));
-    Vec i4;
-    i4[0] = val & 0xffffffff;
-    i4[1] = flag;
-    i4[2] = (val >> 32);
-    i4[3] = flag;
-    asm volatile ("flat_store_dwordx4 %0, %1 sc0 sc1 nt" :: "v"(dst), "v"(i4));
-#else
     union ncclLLFifoLine i4;
     i4.data1 = val & 0xffffffff;
     i4.flag1 = flag;
@@ -277,6 +268,8 @@ private:
     i4.flag2 = flag;
     __builtin_nontemporal_store(i4.v[0], dst->v);
     __builtin_nontemporal_store(i4.v[1], dst->v+1);
+#if defined(__gfx950__)
+    __builtin_amdgcn_fence(__ATOMIC_RELEASE, ""); // flush cache
 #endif
 #else
     asm volatile("st.volatile.global.v4.u32 [%0], {%1,%2,%3,%4};" :: "l"(&dst->i4), "r"((uint32_t)val), "r"(flag), "r"((uint32_t)(val >> 32)), "r"(flag) : "memory");
@@ -351,6 +344,9 @@ private:
       __builtin_nontemporal_store(u4, (uint32_t*)dst);
     else
       __builtin_nontemporal_store(u8, (uint64_t*)dst);
+#if defined(__gfx950__)
+    __builtin_amdgcn_fence(__ATOMIC_RELEASE, ""); // flush cache
+#endif
 #else
     if(sizeof(U) == 1)
       asm volatile("st.volatile.global.b8 [%0],%1;" :: "l"(dst), "r"(u4) : "memory");
@@ -432,11 +428,7 @@ private:
   }
 
   template <int RECV, int SEND, int SrcBuf, int DstBuf>
-#if defined(__gfx950__)
-  __device__ __attribute__((noinline)) void LLGenericOp(intptr_t srcIx, intptr_t dstIx, int nelem, bool postOp) {
-#else
   __device__ void LLGenericOp(intptr_t srcIx, intptr_t dstIx, int nelem, bool postOp) {
-#endif
     constexpr int SRC = SrcBuf != -1 ? 1 : 0;
     constexpr int DST = DstBuf != -1 ? 1 : 0;
     T *srcElts = SrcBuf == -1 ? nullptr : userBufs[SrcBuf] + srcIx;