AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

- LL Protocol: Add missing fences for gfx950, this fixes the hang issue (#1932)

Browse Source 
- Remove asm flat_store_dwordx4, not needed
This commit is contained in:
Larry Meadows
2025-09-25 14:07:07 -07:00
committed by GitHub
parent 01d16d4139
commit cb14fccdcc
1 changed files with 5 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/device/prims_ll.h
+5 -13
View File
@@ -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;