Skip to content

fix(rccl): LL protocol explicit 128-bit vector loads/stores #10670

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
alex-breslow-amd merged 5 commits into from
Sep 26, 2026
Merged

fix(rccl): LL protocol explicit 128-bit vector loads/stores #10670

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
908935b
fix(rccl): use explicit 128-bit vector loads/stores in LL protocol
alex-breslow-amd Aug 24, 2026
4298987
Remove cruft
alex-breslow-amd Aug 25, 2026
808dc4f
Update projects/rccl/src/include/device.h
alex-breslow-amd Aug 25, 2026
eb3a2ab
Assert ncclLLFifoLine stays 16-byte sized and aligned.
alex-breslow-amd Sep 24, 2026
1726a50
Reword LL 128-bit load/store comments to the instruction-selection ra…
alex-breslow-amd Sep 25, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
59 changes: 29 additions & 30 deletions projects/rccl/src/device/prims_ll.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,26 @@
************************************************************************/

#include "rccl_ptr.h"

// Non-temporal 128-bit load for LL communication (FIFO) buffers. Comm buffers
// are uncached; a single explicit vector load avoids relying on the backend to
// fuse two independent 64-bit nontemporal loads into one 16-byte op.
inline __device__ void loadLLLine(const union ncclLLFifoLine* src, union ncclLLFifoLine& dst) {
dst.v4u = __builtin_nontemporal_load((v4u_gptr)&src->v4u);
}

// Plain 128-bit vector store for LL FIFO lines. Comm buffers are uncached; a
// single explicit vector store avoids relying on the backend to fuse two
// independent 64-bit stores into one 16-byte op.
inline __device__ void storeLLLine(union ncclLLFifoLine* dst, const union ncclLLFifoLine& src) {
#if !RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS && (defined(__gfx1200__) || defined(__gfx1201__))
__scoped_atomic_store_n((u64_gptr)dst->v, src.v[0], __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((u64_gptr)dst->v + 1, src.v[1], __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM);
Comment thread
alex-breslow-amd marked this conversation as resolved.
Resolved
#else
*((v4u_gptr)&dst->v4u) = src.v4u;
#endif
}

// UserRegMode is accepted only to match the Primitives primary template (which
// carries it for LL128, see primitives.h/prims_ll128.h). The LL protocol path is
// unchanged from the baseline and ignores it.
Expand Down Expand Up @@ -175,15 +195,9 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
"s_waitcnt vmcnt(0)\n"
: "=v"(i4.i4)
: "v"(&src->i4));
#elif RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
// Comm FIFO buffers are uncached; use non-temporal loads so the flag poll
// never observes a stale cache line (no system-scope cache-bypass needed).
// Holds for the legacy IPC allocator only. See loadLLLineB128.
*((u64_gptr)i4.v) = __builtin_nontemporal_load((u64_gptr)src->v);
*((u64_gptr)i4.v + 1) = __builtin_nontemporal_load((u64_gptr)src->v + 1);
#else
*((u64_gptr)i4.v) = __builtin_nontemporal_load((u64_gptr)src->v);
*((u64_gptr)i4.v + 1) = __builtin_nontemporal_load((u64_gptr)src->v + 1);
// Holds for the legacy IPC allocator only. See loadLLLineB128.
loadLLLine(src, i4);
#endif
if (checkAbort(abort, 1, spins)) break;
} while ((i4.flag1 != flag) || (i4.flag2 != flag));
Expand Down Expand Up @@ -218,14 +232,10 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
"s_waitcnt vmcnt(0)\n"
: "=v"(line[i].i4)
: "v"(&src->i4));
#elif RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
#else
// Comm FIFO buffers are uncached; use non-temporal loads (no bypass).
// Holds for the legacy IPC allocator only. See loadLLLineB128.
line[i].v[0] = __builtin_nontemporal_load((u64_gptr)src->v);
line[i].v[1] = __builtin_nontemporal_load((u64_gptr)src->v + 1);
#else
line[i].v[0] = __builtin_nontemporal_load((u64_gptr)src->v);
line[i].v[1] = __builtin_nontemporal_load((u64_gptr)src->v + 1);
loadLLLine(src, line[i]);
#endif
#else
asm volatile("ld.volatile.global.v4.u32 {%0,%1,%2,%3}, [%4];"
Expand All @@ -250,14 +260,10 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
"s_waitcnt vmcnt(0)\n"
: "=v"(line[i].i4)
: "v"(&src->i4));
#elif RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
#else
// Comm FIFO buffers are uncached; use non-temporal loads (no bypass).
// Holds for the legacy IPC allocator only. See loadLLLineB128.
line[i].v[0] = __builtin_nontemporal_load((u64_gptr)src->v);
line[i].v[1] = __builtin_nontemporal_load((u64_gptr)src->v + 1);
#else
line[i].v[0] = __builtin_nontemporal_load((u64_gptr)src->v);
line[i].v[1] = __builtin_nontemporal_load((u64_gptr)src->v + 1);
loadLLLine(src, line[i]);
#endif
#else
asm volatile("ld.volatile.global.v4.u32 {%0,%1,%2,%3}, [%4];"
Expand All @@ -284,17 +290,10 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
// writer's cache and never reach the peer's poll. One b128 keeps data+flag
// in a single transaction. See RCCL_LL_FIFO_SYS_SCOPE in rccl_ptr.h.
__builtin_amdgcn_global_store_b128((v4u_gptr)dst, i4.v4u, RCCL_SYSTEM_SYNCSCOPE);
#elif RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
*((u64_gptr)dst->v) = *((u64_gptr)i4.v);
*((u64_gptr)dst->v + 1) = *((u64_gptr)i4.v + 1);
#else
#if defined(__gfx1200__) || defined(__gfx1201__)
__scoped_atomic_store_n((u64_gptr)dst->v, i4.v[0], __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
__scoped_atomic_store_n((u64_gptr)dst->v + 1, i4.v[1], __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM);
#else
*((u64_gptr)dst->v) = *((u64_gptr)i4.v);
*((u64_gptr)dst->v + 1) = *((u64_gptr)i4.v + 1);
#endif
// Comm FIFO buffers are uncached. storeLLLine emits a plain 128-bit vector
// store, or a system-scope release pair on gfx1200/gfx1201.
storeLLLine(dst, i4);
#endif
#if defined(__gfx950__) && ROCM_VERSION < 70002
__builtin_amdgcn_fence(
Expand Down
10 changes: 5 additions & 5 deletions projects/rccl/src/include/device.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,9 +121,8 @@ struct ncclDevRedOpFull {
uint64_t scalarArg;
};

#ifdef __HIP_DEVICE_COMPILE__
#include "rccl_ptr.h"
#endif

#include "nccl_device/rccl_ptr.h"

union ncclLLFifoLine {
/* Flags have to be *after* data, because otherwise, an incomplete receive
Expand All @@ -138,11 +137,12 @@ union ncclLLFifoLine {
};
uint64_t v[2];
int4 i4;
#if defined(__HIP_DEVICE_COMPILE__) && RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
v4u v4u; /* same layout as data1,flag1,data2,flag2 for b128 load/store */
#endif
};

static_assert(sizeof(union ncclLLFifoLine) == 16, "ncclLLFifoLine must stay 16 bytes");
static_assert(alignof(union ncclLLFifoLine) == 16, "ncclLLFifoLine must stay 16-byte aligned");

#if __HIP_DEVICE_COMPILE__
#if defined(__GFX9__)
#define WARP_SIZE 64
Expand Down
Loading