Repository navigation
fix(rccl): drop the hipHostMallocUncached fence from LL store<> - #13451
Merged
Merged
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
There was a problem hiding this comment.
🟢 Approval recommended
The focused implementation preserves intended memory scope, removes the identified bottleneck, and is supported by extensive performance and correctness validation.
0 open findings
What changed in this PR
Ensures LL user-buffer stores retain system-scope semantics on gfx942/gfx950 without b128 builtins, eliminating unnecessary gfx950 release fences.
Changes:
- Uses scoped atomic stores based on GPU architecture.
- Removes the now-unreachable gfx950 fallback fence.
- Documents the latency fix.
| File | Description |
|---|---|
projects/rccl/src/device/prims_ll.h |
Corrects LL user-output store selection and removes redundant fencing. |
projects/rccl/CHANGELOG.md |
Records the performance regression fix. |
🧠 Review effort: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
storeData has called store<> since #9336. When RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS is 0 (compiler without __builtin_amdgcn_global_load_b128, e.g. ROCm 7.0.1, or DWORDX4_INTRINSICS=OFF), store<> falls back to a non-temporal store followed, on gfx950 with ROCM_VERSION < 70002, by a release fence: buffer_wbl2 sc0 sc1 plus s_waitcnt vmcnt(0) lgkmcnt(0) after every user-buffer store. The fence works around the ROCr hipHostMallocUncached bug for the comm FIFO, which storeLL still covers. store<> only writes user buffers, had no callers before #9336, and its builtin path has not carried the fence since #3812. Only gfx950 builds against ROCm < 7.0.2 without the builtins change; every other build preprocesses to the same code. JIRA ID : ROCM-31938 Co-authored-by: Cursor <cursoragent@cursor.com>
i-kosarev
changed the title
i-kosarev
force-pushed
the
users/ilkosare/rocm31938-bcast
branch
from
October 14, 2026 13:03
6a7c228 to
d3794a3
Compare
i-kosarev
requested review from
alex-breslow-amd and
wenkaidu
and
a balanced review from Copilot
October 14, 2026 14:37
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Review fix: the comment left in place of the fence described code that is no longer there. store<> now says up front that it writes user buffers only and that comm FIFO stores must go through storeLL, which keeps the gfx950 hipHostMallocUncached fence for ROCm < 7.0.2. Comment only. JIRA ID : ROCM-31938 Co-authored-by: Cursor <cursoragent@cursor.com>
amd-arozanov
approved these changes
Oct 15, 2026
i-kosarev
added a commit
that referenced
this pull request
Oct 15, 2026
1 task done
arvindcheru
pushed a commit
that referenced
this pull request
Oct 16, 2026
…er peer (#13468) ## Motivation Commit 242e0a9 (#11635) raised `p2pnChannelsPerPeer` from 1 to 2 on AINIC for jobs of up to 8 nodes, so that a single peer can saturate a NIC. That helps point-to-point traffic, but `ncclScatter` regressed by 15 to 21% from 32 MB upward on multi-node AINIC jobs compared with 7.14. Scatter already drives all `nRanks - 1` peers from one root, so splitting each peer into two parts adds no bandwidth and costs some. This PR restores scatter performance without changing the channel count for the other collectives. ## Technical Details - New helper `rcclP2pTaskParts(comm, collAPI)` returns how many of the `p2pnChannelsPerPeer` parts a P2P task uses and connects. For `ncclFuncScatter` on multi-node AINIC it returns the largest power of two not above `p2pnChannels / (nRanks - 1)`, capped at `p2pnChannelsPerPeer` and at least 1, so the root's flows still fit the P2P channel pool. A user-set `NCCL_NCHANNELS_PER_NET_PEER` is kept, as in `ncclTopoComputeP2pChannels`. Otherwise it returns `p2pnChannelsPerPeer`. - The decision depends only on `collAPI` and comm-wide sizes, so sender and receiver always agree on the number of parts. - Limited to AINIC: it is the only NIC path that raises `p2pnChannelsPerPeer` back to 2 on multi-node (#11635), and the only one measured. On other NICs the halving loop in `ncclTopoComputeP2pChannels` already reduces the parts per peer when the peers do not fit the pool. - Used in three places: - channel count selection (`rcclP2pPolicyChannels`), also with P2P disabled, since the limit is about network flows; - the connector checks and buffer registration in `addP2pToPlan`; - preconnect (`p2pTaskAppend`, `postTuneP2pRecordPreconnect`). - NIC buffer registration (`postTuneP2pRegisterBuffer`) keeps `p2pnChannelsPerPeer`: Scatter tasks have `allowUB` off and never reach it. - When only some parts are connected, `sendSeen` / `recvSeen` stay clear, so a later task to the same peer that uses all parts still connects the rest. - For every other collective and for send/recv the helper returns `p2pnChannelsPerPeer`, which equals `nChannelsMax` in `addP2pToPlan`, so their code path takes the same decisions as before. - Host microtests: - `EnqueueMicrotest.P2pTaskParts_*`: the part rule (pools of 2, 3, 4, 8 and 16 nodes, the pool boundary, power-of-two rounding, the one-part floor, non-AINIC, single node, user-set count); - `EnqueueMicrotest.P2pPolicyChannels_ScatterBelowPerPeerReturnsItsParts`: the channel count; - `TaskPostTuningMicrotest.P2pRecordPreconnect_ScatterWithFewerParts_ClaimsItsPartsAndLeavesTheRestForLater`: the partial-to-full preconnect on both sides. Limiting only the channel count is not enough. When scatter sends on part 0 but connects both parts itself, it stays at or below the two-channel speed in most runs (4 of 5 runs at about 110 GB/s instead of 137 GB/s). Not connecting part 1 makes the result stable. Part 1 connected later by another collective on the same communicator (alltoall) does not slow scatter down. The exact runtime cause of the slowdown is not identified yet: the connection setup logs of fast and slow runs are identical. With `NCCL_PXN_DISABLE=1` there is no regression, because the root is limited by its own NIC. ## Issue Tracking JIRA ID : ROCM-31938 ## Test Plan - 4 nodes x 8 MI355X (32 ranks), AINIC (`NCCL_NET=ROCM-IB`, `RCCL_AINIC_ROCE=1`), PXN on. - rccl-tests with data check on (`-c 1`), `-n 20 -w 5`, 5 cycles per run. - Three builds from the same base, run back to back in every repeat: - revert: develop without the AINIC `p2pnChannelsPerPeer` raise from 242e0a9 (matches 7.14); - develop; - this PR. - All P2P collectives (scatter, gather, alltoall, alltoallv, all_gather, reduce_scatter) up to 16 GB, 2 repeats. - All collectives (also sendrecv, all_reduce, broadcast, reduce) up to 2 GB, 3 repeats. - Scatter on one communicator before and after alltoall and after ring sendrecv, with every received byte checked. - Scatter with `NCCL_PXN_DISABLE=1` and on 2 nodes. - Scatter 16M–1G, root 0, on 2, 3, 4, 5 and 8 nodes, interleaved with develop and the first version of this PR (one part per peer whenever the parts did not fit); all roots on 3 and 4 nodes (from #13528). - Host microtests `EnqueueMicrotest.*` and `TaskPostTuningMicrotest.*`. ## Test Result Scatter, 4 nodes, 32 ranks, median time over 10 samples. Measured with the first version of this PR; on 4 nodes the current rule picks the same single part. | Size | revert (us) | develop (us) | this PR (us) | revert busbw (GB/s) | develop busbw (GB/s) | this PR busbw (GB/s) | develop / revert | this PR / revert | |---:|---:|---:|---:|---:|---:|---:|---:|---:| | 8M | 48.9 | 49.9 | 48.6 | 166.2 | 162.9 | 167.2 | 1.02 | 0.99 | | 16M | 69.9 | 73.3 | 70.1 | 232.5 | 221.7 | 231.9 | 1.05 | 1.00 | | 32M | 109.2 | 132.2 | 111.0 | 297.7 | 245.9 | 292.8 | 1.21 | 1.02 | | 64M | 190.4 | 228.6 | 190.7 | 341.4 | 284.4 | 340.9 | 1.20 | 1.00 | | 128M | 346.9 | 413.8 | 347.5 | 374.8 | 314.2 | 374.2 | 1.19 | 1.00 | | 256M | 655.5 | 768.9 | 655.3 | 396.7 | 338.2 | 396.8 | 1.17 | 1.00 | | 512M | 1281.6 | 1490.8 | 1278.1 | 405.8 | 348.9 | 406.9 | 1.16 | 1.00 | | 1G | 2537.6 | 2951.9 | 2534.4 | 409.9 | 352.4 | 410.4 | 1.16 | 1.00 | | 2G | 5048.6 | 5865.6 | 5050.6 | 412.1 | 354.7 | 411.9 | 1.16 | 1.00 | | 4G | 10082.9 | 11698.2 | 10084.8 | 412.7 | 355.7 | 412.6 | 1.16 | 1.00 | | 8G | 20167.1 | 23342.6 | 20175.5 | 412.6 | 356.5 | 412.5 | 1.16 | 1.00 | | 16G | 40433.2 | 46671.4 | 40393.4 | 411.6 | 356.6 | 412.0 | 1.15 | 1.00 | Scatter by node count, root 0, time ratio vs develop over 16M–1G (from #13528, merged here): | Nodes | Parts per peer (develop / first version / this PR) | first version | this PR | | --- | --- | --- | --- | | 2 | 2 / 2 / 2 | 0.99–1.01 | 0.96–1.00 | | 3 | 4 / 1 / 2 | 1.04–1.13 | 0.85–0.87 | | 4 | 2 / 1 / 1 | 0.80–0.86 | 0.79–0.86 | | 5 | 2 / 1 / 1 | 0.92–1.11 | 0.91–1.08 | | 8 | 2 / 1 / 1 | 0.97–1.02 | 0.94–1.00 | - Scatter matches revert from 8 MB to 16 GB. Develop is 15 to 21% slower from 32 MB up. - 3 nodes: the first version dropped 4 parts to 1 and left two-thirds of the 64-channel pool idle; this PR keeps the 2 parts that fit. All roots on 3 nodes: median 0.82–0.90 vs develop (first version: 0.99–1.10). All roots on 4 nodes: both at 0.84–0.90. - Scatter after alltoall on the same communicator stays at revert level (1 GB: about 2530 us, develop about 2980 us). - With root != 0, scatter shows run-to-run bimodality (about 2.5 vs about 3.5 ms at 1 GB), as in the revert build (pre-242e0a93afa, one part per peer); develop, with two parts per peer, does not show it. The median over all roots is 0.82–0.89 of develop, and root 0 is stable at revert level in every run. - gather, all_gather, reduce_scatter, all_reduce, broadcast, reduce, sendrecv: no difference from develop beyond run-to-run noise. Outliers above 5% only at latency-bound sizes from 1 KB to 1 MB (a few us at a 20 to 30 us spread), also in collectives that do not go through the changed code. - alltoall and alltoallv show the same difference against revert in develop and in this PR (alltoall: develop 1.05, this PR 1.02 to 1.045 at 256 MB to 16 GB). This PR does not add to it. - 2 nodes: no regression in develop, this PR is unchanged. - `NCCL_PXN_DISABLE=1`: all builds at about 64 GB/s at 256 MB, no regression. - Joint validation on ROCm 10.2 (branch ROCM-31938-validation = develop + #13451 + this PR + #13536, 4 nodes, median of 10 from two launches in opposite build order): scatter from 16 MB 0.78–0.86 of develop; against 7.2.3 16 MB 1.13, 32 MB 1.07, from 128 MB 0.95–0.98. The other collectives match develop within noise; on alltoallv two copies of the same develop build differ by up to 6% (A/A run), and validation is 0.98–1.04 of develop over 20 measurements. - Correctness: `#wrong = 0` and no out-of-bounds values in every run. - Host microtests `EnqueueMicrotest.*` and `TaskPostTuningMicrotest.*` pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Ilia Kosarev <Ilia.Kosarev@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
1 task
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.
Motivation
Since #9336, the LL
storeDatawrites the user output buffer throughstore<>. Without the b128 builtins (RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS == 0),store<>uses a non-temporal store. On gfx950 withROCM_VERSION < 70002it also issues a release fence (buffer_wbl2 sc0 sc1+s_waitcnt vmcnt(0) lgkmcnt(0)) after every store. This makes 4 KB–64 KBncclBroadcaston MI355X about 30% slower in RCCL built against ROCm 7.0.1.Technical Details
ROCM_VERSION < 70002fence from thestore<>fallback.hipHostMallocUncachedworkaround for the comm FIFO, andstoreLLkeeps it.store<>only writes user buffers, and its builtin path has had no fence since Fully enable cache bypass access for Simple/LL/LL128 #3812.Issue Tracking
JIRA ID: ROCM-31938
Test Plan
4 nodes x 8 MI355X, RCCL built against ROCm 7.0.1:
broadcast_perf1K–256K, interleaved with RCCL 2.27.7 and develop.all_reduce_perfandreduce_perf1K–4M.NCCL_PROTO=LL -d all -u 1, 8 B–1 MB, for broadcast, all_reduce, all_gather, reduce_scatter and reduce.Test Result
broadcast_perfout-of-place time, ratio vs RCCL 2.27.7:all_reduceandreduce: no change.#wrong= 0 for every collective, data type and size.Submission Checklist