Repository navigation
fix(rccl): use MI450 cooperative 128-bit builtins and explicit vector loads/stores in LL128 - #10666
Open
alex-breslow-amd wants to merge 9 commits into
Open
alex-breslow-amd wants to merge 9 commits into
alex-breslow-amd wants to merge 9 commits into
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Contributor
There was a problem hiding this comment.
Pull request overview
This PR updates RCCL LL128 128-bit memory operations for gfx1250/MI450 and fallback architectures.
Changes:
- Adds vector pointer types and device synchronization scope definitions.
- Adds gfx1250 cooperative atomic load/store paths.
- Replaces paired 64-bit accesses with explicit vector operations.
Reviewed changes
Copilot reviewed 3 out of 3 changed files in this pull request and generated 3 comments.
| File | Summary | Findings |
|---|---|---|
projects/rccl/src/include/nccl_device/rccl_ptr.h |
Adds vector types and sync-scope definitions. | None. |
projects/rccl/src/device/prims_ll128.h |
Uses explicit vector non-temporal loads. | None. |
projects/rccl/src/device/op128.h |
Implements cooperative and vector-based 128-bit operations. | Two critical findings: cooperative load/store calls may lack full 8-lane participation. One moderate finding: the fallback load uses device scope instead of the required system scope. |
Suppressed comments (1)
projects/rccl/src/device/op128.h:22
- This adds the gfx1250 cooperative path, but the current device tests do not execute an LL128 collective with registered user buffers:
TestOp128only round-trips one directload128/store128pair, andLL128RegVariantKernel_testchecks function-table IDs rather than memory behavior. Add a gfx1250 multi-rank registered-buffer LL128 data round-trip (including the cooperative load/store path) before relying on this correctness-critical change.
#if __has_builtin(__builtin_amdgcn_cooperative_atomic_load_8x16B) && defined(__gfx1250__)
u.vi = __builtin_amdgcn_cooperative_atomic_load_8x16B((v4i_gptr)ptr, __ATOMIC_RELAXED, RCCL_DEVICE_SYNCSCOPE);
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| BytePack<16> ans; \ | ||
| ans.u64[0] = *((uint64_t*)addr); \ | ||
| ans.u64[1] = *((uint64_t*)addr + 1); \ | ||
| ans.v4u = *((v4u_gptr)addr); \ |
Contributor
Author
There was a problem hiding this comment.
This change is probably a performance win. I need to check if it now outputs global instead of flat instructions. Previously, that helped a fair bit. The compiler can often lose track of whether a pointer is to global memory or LDS. In that ambiguous case, it emits a flat vector memory instruction, which works on both. However, that can incur a performance hit since the hardware needs to do extra work. I'll have to take a look at the assembly to see if even this is working. Previously, I had to get the address of ans.v4u, cast it to an address_space(1) pointer and then dereference it again. Gross, but you get the idea.
1 of 4 tasks
Contributor
Author
|
Stack: #10670 is stacked on this PR.
Merge order: #10666 → then #10670 (GitHub will retarget #10670 to `develop` after this lands). |
3 tasks
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
ARGUS
Findings
1 finding(s) that fall outside the diff and could not be anchored to a line
projects/rccl/CHANGELOG.md:24(Nit) No CHANGELOG entry for the gfx1250 LL128 load/store change; the unreleased 2.30.7 section has none and the file carries device-internal entries at this granularity.
Bottom line: Worth settling before approval: the new gfx1250 guard at op128.h:21 skips DWORDX4_INTRINSICS_FORCE_OFF, so -DDWORDX4_INTRINSICS=OFF no longer disables 128-bit builtins there. Everything else I found is Low or Nit.
Comment @argus review to re-run, or @argus re-review after pushing.
alex-breslow-amd
force-pushed
the
users/abreslow/ll128_cleanup
branch
from
October 1, 2026 00:04
644ff87 to
b6b6819
Compare
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: -DDWORDX4_INTRINSICS=OFF still does not disable the new cooperative branch at op128.h:21, because COOPERATIVE_ATOMIC_INTRINSICS_FORCE_OFF (rccl_ptr.h:56) is defined nowhere in the build.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
|
AI review of this PR (particularly the review of cooperative atomics) says that the cooperative atomics aren't needed where they are used, and are not used where they are needed. They should be used in recvReduceSendCopy (in prims_ll128.h) in place of load128NT and store128Fifo. I'll admit I haven't studied the code yet myself, but the analysis seems correct. |
mustafabar
reviewed
Oct 1, 2026
Contributor
Author
Yup, I messed up. :) |
Contributor
There was a problem hiding this comment.
ARGUS
Findings
1 finding(s) that could not be posted as inline comments
projects/rccl/src/device/prims_ll128.h:598(Low) skip_fence's stated justification names global_store_b128, which is no longer the FIFO store emitted on gfx1250skip_fenceat prims_ll128.h:596 is gated onRCCL_HAVE_GLOBAL_DWORDX4_BUILTINS, and the comment at prims_ll128.h:598 justifies the cheap fence by the stores being system-scopeglobal_store_b128, but on gfx1250 the FIFO store is nowcooperative_atomic_store_8x16B. Those lines sit outside this diff, so worth a follow-up rather than a change here. Please check.
Bottom line: Worth settling before approval: the inline asm in load128Fifo and store128Fifo (prims_ll128.h:48, :74) has write-only outputs under an empty template, so neither one actually moves its value.
Comment @argus review to re-run, or @argus re-review after pushing.
| typedef __attribute__((address_space(1))) v4i* v4i_gptr; | ||
|
|
||
| // One 16-byte pun for those two vector types and a pair of uint64 words. | ||
| union alignas(16) rcclB128 { |
Contributor
There was a problem hiding this comment.
Low
rcclB128 has no uint32_t u32[4], so the 32-bit-word puns at prims_ll.h:95, algorithms/dda/device/CollCommon.h:261 and nccl_device/impl/ll_a2a__funcs.h:18 cannot adopt it. Those files are outside this PR, so just noting it once: adding the member later would let the consolidation finish.
Contributor
Author
There was a problem hiding this comment.
Agreed. I'm leaving this for the follow-up that consolidates the 128-bit puns (prims_ll.h, DDA) onto rcclB128, to keep this PR scoped to LL128.
Contributor
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
| rcclB128 u; | ||
| // See rcclIdentity in rccl_ptr.h. | ||
| u.vi = rcclIdentity( | ||
| __builtin_amdgcn_cooperative_atomic_load_8x16B((v4i_gptr)ptr, __ATOMIC_RELAXED, RCCL_SYSTEM_SYNCSCOPE)); |
Contributor
Author
There was a problem hiding this comment.
I tried a bunch of simpler things here, but defining this rcclIdentity function was the only thing that worked and didn't emit additional instructions.
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: no test distinguishes the gfx1250 cooperative FIFO load/store from the system-scope b128 fallback, so a revert of load128Fifo/store128Fifo would pass unnoticed.
Comment @argus review to re-run, or @argus re-review after pushing.
| rcclB128 u; | ||
| // See rcclIdentity in rccl_ptr.h. | ||
| u.vi = rcclIdentity( | ||
| __builtin_amdgcn_cooperative_atomic_load_8x16B((v4i_gptr)ptr, __ATOMIC_RELAXED, RCCL_SYSTEM_SYNCSCOPE)); |
Contributor
There was a problem hiding this comment.
Medium
Medium
load128Fifo/store128Fifo move gfx1250 LL128 from the system-scope b128 to the cooperative builtin. LlFifoSysScope.SiblingBroadcastSlotReuse_LL128 passes on either branch, and it skips outright without DPX siblings (LlFifoSysScopeTests.cpp:238). Could a test pin the selected path, for example a single-GPU 8-lane cooperative round-trip kernel built without DWORDX4_INTRINSICS_FORCE_OFF?
Contributor
Author
There was a problem hiding this comment.
Added in 2dac0b5:
LlFifoSysScopeDeviceTest.Gfx1250SelectsCooperativechecks in the device pass thatRCCL_HAVE_COOPERATIVE_ATOMIC_BUILTINSis 1 on gfx1250 unless the builtins are forced off, and 0 elsewhere.LlFifoSysScopeDeviceTest.CooperativeLineRoundtripis a single-GPU cooperative store/load round trip over whole waves (each 8-lane group is one 128-byte line). It also reloads and masks the lines in a second kernel, which crashes ROCm 10.1 instruction selection ifrcclIdentityis removed.
Both compile for gfx942 and gfx1250. They have not run on MI450 yet, so I'm leaving this open until they do.
Contributor
There was a problem hiding this comment.
Resolved by 813afbc39a.
| // to a copy, but the SelectionDAG cannot see through it. | ||
| #if ROCM_VERSION < 100200 | ||
| extern "C" __device__ __attribute__((const, noconvergent)) v4i rcclSoftWqmV4i(v4i) | ||
| __asm("llvm.amdgcn.softwqm.v4i32"); |
Contributor
Author
There was a problem hiding this comment.
Ordinarily this would be unsafe. However, we've clamped the compiler version to one where this LLVM IR exists.
Contributor
Author
|
Re the outside-diff |
Contributor
|
@cppccppccppc @sunguobao PTAL |
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: the new cooperative tests build in rccl-UnitTestsFixtures under DWORDX4_INTRINSICS_FORCE_OFF, so they never compile the 8x16B builtins or rcclIdentity.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
pr-rocm-systems-10666-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. |
Contributor
pr-rocm-systems-10666-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. |
…d gate it per arch Follow-up to #14128 (commit 1, c34bb45, "perf(rccl): keep LL128 lane state in registers"). #14128 fixes a real LL128 bottleneck: the Primitives object lives in scratch, so the slice loop reloaded the lane id and peer 0's FIFO pointer and flag after the remote FIFO stores, and on gfx9 the vmcnt(0) guarding that reload also waits for those stores. This change keeps the fix and changes how it is expressed: - Keep NCCL's names. wid and flagThread stay at all 22 use sites in prims_ll128.h instead of being renamed to laneId()/isFlagLane(). When the feature is enabled they are [[no_unique_address]] empty members whose conversion operators read __lane_id(), so they take no storage and the generated code is the same as #14128's. - Gate it per arch at compile time. RCCL_LL128_HW_LANE_ENABLE is defined for __gfx1250__/__gfx950__. Every other arch keeps the original stored members and the per-slice reads of peer 0's FIFO pointer and flag. Enabling another arch is a one-line change. - Revert the rename in tdm/ll128Tdm.h. Why keep the names: RCCL syncs prims_ll128.h with NCCL. NCCL's prims_ll128.h (2.30 and 2.32) uses the same wid and flagThread members on 28 lines, and with this change 18 of RCCL's lines that use them are identical to NCCL's. With the rename, every NCCL sync conflicts on those lines, and any new NCCL code that uses wid or flagThread stops compiling (the members no longer exist) and has to be rewritten on each sync. Keeping the names leaves those lines identical to NCCL, and new NCCL code compiles unchanged and picks up the hardware lane. Why opt in per arch: the change is device code in the LL128 slice loop of every collective (and send/recv where it uses LL128) on every arch. A compile-time gate lets each arch adopt it after validation, with no runtime check in the hot loop and no duplicate kernels. Disabled, the code is the code from before #14128. Correctness: - The hardware lane equals tid % WARP_SIZE only if every LL128 Primitives group starts on a warp boundary. Every construction site meets this: the LL128 tree split rounds to whole warps, the batch dispatcher uses nWarps * WARP_SIZE threads, the warp-speed path passes tid % WARP_SIZE, other offset-tid constructions are Simple-only, and WARP_SIZE matches the hardware wave (64 on GFX9, 32 otherwise). The requirement is documented next to the members. - Peer 0's step only advances after the slice loop, and waitSend runs before the hoisted values are read, so computing them once per op gives the same values. - wid converts implicitly to int and flagThread to bool, which covers every current use (subscripts, arithmetic, comparisons, ll128SlotIx<...>(g, wid), !, &&, the ternary). Code that needs a real int (auto, varargs, template deduction) should write int(wid). Validation (on a tree carrying #14128 commit 1 backported to an older drop; per-function machine-code comparison of the device code objects): - gfx950 (enabled): 9146/9146 functions byte-identical to #14128. - gfx942 (disabled): 9279/9279 functions byte-identical to before #14128. gfx950 with its guard removed: 9150/9150 byte-identical to before #14128, so the disabled path is the pre-#14128 code. - Hardware, gfx950 enabled, same rccl-tests sweeps as #14128 (bf16; AllReduce, ReduceScatter, AllGather, AllToAll; 3 jobs per scale; same flags and placement): MI355X (Pollara) 1/2/4/8 nodes and MI350X (Thor2) 1/2/4/8/16 nodes, median bus bandwidth within +/-0.5% of #14128 for every collective and scale; 0 wrong values. - gfx1250: the technique (empty members, implicit conversions in every usage pattern above, and a static_assert that the members take no storage) compiles for gfx1250 with TheRock clang. Not run on gfx1250 hardware. For reference, #14128 itself on gfx950 (MI350X and MI355X, 1-16 nodes, bf16 rccl-tests, compared with the same build without it): scratch loads inside loops -16%, vmcnt(0) waits inside loops -6%, instructions -4%; bus bandwidth at LL128 sizes on 2-16 nodes AllReduce +25% to +29%, ReduceScatter +17% to +21%, AllGather +5% to +10%; other sizes within noise; 0 wrong values in 23,280 result rows.
alex-breslow-amd
commented
Nov 3, 2026
Comment on lines
1453
to
1497
|
|
||
| # Print whether gfx1250 cooperative 128-bit atomics are actually selected, per | ||
| # GPU target, on every build. A #pragma message in rccl_ptr.h would not show up: | ||
| # the device linker compiles with -w, which drops it. This syntax-only compile | ||
| # does not. | ||
| if(REPORT_COOPERATIVE_ATOMICS) | ||
| set(_rccl_coop_log_src "${CMAKE_CURRENT_BINARY_DIR}/rccl_coop_atomic_log.hip") | ||
| file(WRITE "${_rccl_coop_log_src}" [=[ | ||
| #ifdef __HIP_DEVICE_COMPILE__ | ||
| #include "rccl_ptr.h" | ||
| #if RCCL_HAVE_COOPERATIVE_ATOMIC_BUILTINS | ||
| #pragma message("RCCL cooperative atomic builtins: on (" __amdgcn_processor__ ")") | ||
| #else | ||
| #pragma message("RCCL cooperative atomic builtins: off (" __amdgcn_processor__ ")") | ||
| #endif | ||
| #endif | ||
| __global__ void rcclCoopAtomicLog() {} | ||
| ]=]) | ||
| set(_rccl_coop_log_flags | ||
| -x hip -std=c++17 -fsyntax-only -Wno-unused-command-line-argument | ||
| -I${RCCL_SOURCE_DIR}/src/include/nccl_device | ||
| ) | ||
| if(NOT DWORDX4_INTRINSICS) | ||
| list(APPEND _rccl_coop_log_flags -DDWORDX4_INTRINSICS_FORCE_OFF) | ||
| endif() | ||
| # Target features (e.g. :xnack+) do not change the result. | ||
| set(_rccl_coop_archs "") | ||
| foreach(_rccl_coop_arch_raw ${GPU_TARGETS}) | ||
| string(REGEX REPLACE ":.*" "" _rccl_coop_arch "${_rccl_coop_arch_raw}") | ||
| list(APPEND _rccl_coop_archs "${_rccl_coop_arch}") | ||
| endforeach() | ||
| list(REMOVE_DUPLICATES _rccl_coop_archs) | ||
| set(_rccl_coop_log_cmds "") | ||
| foreach(_rccl_coop_arch ${_rccl_coop_archs}) | ||
| list(APPEND _rccl_coop_log_cmds | ||
| COMMAND ${CMAKE_CXX_COMPILER} --offload-arch=${_rccl_coop_arch} ${_rccl_coop_log_flags} ${_rccl_coop_log_src}) | ||
| endforeach() | ||
| add_custom_target(rccl_log_cooperative_atomics ALL | ||
| ${_rccl_coop_log_cmds} | ||
| COMMENT "RCCL cooperative atomic builtins" | ||
| VERBATIM | ||
| ) | ||
| endif() | ||
| ## Set RCCL linked library directories | ||
| target_link_directories(rccl PRIVATE ${SMI_LIB_DIR}) |
Contributor
Author
There was a problem hiding this comment.
I think we can remove this check.
Co-authored-by: Cursor <cursoragent@cursor.com>
ROCm 10.1's compiler lowers __builtin_amdgcn_cooperative_atomic_load_8x16B to an i128 ATOMIC_LOAD whose computed known bits have the wrong width. Masks on the result crash instruction selection (APInt::andAssignSlowCase) or silently miscompile. Pass the cooperative load result through rcclIdentity, which on ROCm < 10.2 wraps llvm.amdgcn.softwqm: it lowers to a copy outside WQM, but the SelectionDAG cannot see through it. On ROCm >= 10.2 it is a plain identity. The store side never needed the workaround, so store128Fifo fills the union directly. Co-authored-by: Cursor <cursoragent@cursor.com>
- store128Fifo: delegate the system-scope non-cooperative path to store128. - Update the skip_fence comment to cover the gfx1250 cooperative store. - Make the cooperative-atomic configure report opt-in (-DREPORT_COOPERATIVE_ATOMICS=ON), strip target features and dedupe GPU_TARGETS, honor DWORDX4_INTRINSICS=OFF, and rerun on every build. Co-authored-by: Cursor <cursoragent@cursor.com>
… path - ddaLL128LoadAtomic16B passes the cooperative 8x16B load result through rcclIdentity, like load128Fifo. - LlFifoSysScopeDeviceTest.Gfx1250SelectsCooperative checks that gfx1250 selects the cooperative builtins unless they are forced off. - LlFifoSysScopeDeviceTest.CooperativeLineRoundtrip runs a single-GPU cooperative store/load round trip over whole waves and masks a reloaded line; without rcclIdentity that mask crashes ROCm 10.1 instruction selection. Co-authored-by: Cursor <cursoragent@cursor.com>
…ction rccl-UnitTestsFixtures defines DWORDX4_INTRINSICS_FORCE_OFF, which also forces the cooperative builtins off, so the cooperative tests compiled only their fallback branch and the guard checks compared 0 with 0. Undefine it in this file; its kernels use b128 and cooperative atomics only on gfx1250, so the gfx942 fault behind the define does not apply. Co-authored-by: Cursor <cursoragent@cursor.com>
ROCm 10.2 pre-release compilers (TheRock 10.2.0a20261013) still crash in instruction selection (APInt::andAssignSlowCase) on the raw cooperative_atomic_load_8x16B result, in the fp8 prod Tree LL128 AllReduce kernels. Drop the ROCM_VERSION < 100200 gate so the softwqm identity is always applied on gfx1250. Co-authored-by: Cursor <cursoragent@cursor.com>
alex-breslow-amd
force-pushed
the
users/abreslow/ll128_cleanup
branch
from
November 3, 2026 03:31
b1808a4 to
539fa50
Compare
Drop the REPORT_COOPERATIVE_ATOMICS option and the syntax-only compile that printed, per GPU target, whether the gfx1250 cooperative 128-bit atomics are selected. LlFifoSysScopeDeviceTest.Gfx1250SelectsCooperative already checks the selection on the device. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
pr-rocm-systems-10666-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. pr-rocm-systems-10666-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-10666-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
__builtin_amdgcn_cooperative_atomic_load_8x16B/_store_8x16B) at system scope, so each 8-lane group moves one non-tearing 128-byte line. They live in the newload128Fifo/store128Fifohelpers inprims_ll128.h, whichrecvReduceSendCopyexecutes on every lane.load128,store128,load128NT) stay per-thread, because their callers skip flag lanes and do not run whole groups. They now issue one explicit 16-byte vector operation instead of two 8-byte operations that relied on the compiler fusing them.load128uses system scope.rccl_ptr.haddsv4i/v4i_gptr(the cooperative builtins take signed vectors), a sharedrcclB128pun,RCCL_HAVE_COOPERATIVE_ATOMIC_BUILTINS, andrcclIdentity.-DDWORDX4_INTRINSICS=OFFalso disables the cooperative builtins. Without them, the FIFO falls back to the system-scopeglobal_load_b128/global_store_b128where available, otherwise the previous nontemporal load and plain store.-DREPORT_COOPERATIVE_ATOMICS=ON(default OFF) adds a build target that prints, per GPU target, whether the cooperative builtins are selected.ROCm 10.1 compiler workaround (
rcclIdentity)The ROCm 10.1 LLVM lowers
cooperative_atomic_load_8x16Bto an i128ATOMIC_LOADwhose computed known bits have the wrong width. Masking the result crashes instruction selection (APInt::andAssignSlowCase) or silently miscompiles the mask. Upstream LLVM fixes this (memVT = MVT::getVT(CI.getType())).Every cooperative load result, including DDA's
ddaLL128LoadAtomic16B, passes throughrcclIdentity. OnROCM_VERSION < 100200it callsllvm.amdgcn.softwqmthrough an asm label (no inline asm). Outside WQM that lowers to a copy, so it costs no instructions, but the SelectionDAG cannot see through it. On ROCm 10.2 and later it is a plain identity.JIRA ID: ROCM-31751
Test plan
-O3) and Debug (-O1) builds of all device kernels for gfx1250 with ROCm 10.1-DCOOPERATIVE_ATOMIC_INTRINSICS_FORCE_OFFcompiles and still emits system-scopeglobal_store_b128for the FIFO-DREPORT_COOPERATIVE_ATOMICS=ONreports on for gfx1250 and off for gfx942;gfx942:xnack+;gfx942:xnack-produce one checkLlFifoSysScopeDeviceTest.Gfx1250SelectsCooperativeandCooperativeLineRoundtripcompile for gfx942 and gfx1250;CooperativeLineRoundtripfails to compile on ROCm 10.1 ifrcclIdentityis removedLlFifoSysScopeDeviceTest.*on MI450LlFifoSysScope.SiblingBroadcastSlotReuse_LL128on DPX siblings