Tuning Infra Improvement by isaki001 · Pull Request #12386 · ROCm/rocm-systems

isaki001 · 2026-09-28T17:25:40Z

Motivation

Enable run-time thresholding of Symmetric. Improve run-time threshold for CE.

Technical Details

Env var rename: RCCL_CE_AR_MAX_MSG_BYTES → RCCL_CE_AR_2SHOT_MAX_BYTES.
The name now accurately reflects that this ceiling applies specifically to the two-shot (non-registered) AllReduce path, not
to all CE AllReduce variants.

New runtime threshold overrides (all default to -1 = use arch table):

RCCL_CE_COLL_MAX_BYTES — upper bound for non-AR CE collectives (AllGather, AllToAll)
RCCL_CE_COLL_MIN_BYTES — lower bound for non-AR CE collectives
RCCL_CE_AR_2SHOT_MIN_BYTES — lower bound for CE AllReduce two-shot window
RCCL_SYM_K_MAX_BYTES — upper bound for symmetric kernel selection
RCCL_SYM_K_MIN_BYTES — lower bound for symmetric kernel selection

These are wired into rcclCeRegMaxTab, rcclCeNonRegMaxTab, rcclSymMaxR2CapTab, rcclSymMinR2CapTab, and rcclCeNonRegMinTab respectively.

DDA LL128: Removed rcclParamDdaLL() == 0 guards inside ddaLL128ArOneShotEligible

and ddaLL128ArTwoShotEligible — these conditions are already checked by rcclDdaLL128ArEligible before dispatching into either function.

AllGather dispatch non-registered CE branch fix

Before this PR the AllGather CE-scratch branch fired when either
rcclForceCeScratchEnabled was true or the message fit the size window, making the two conditions alternatives:

Now RCCL_FORCE_CE is required for enablement and cannot be bypassed by tuning thresholds
CTA Policy Zero can bypass thresholds

AllToAll CE registered guard: Added a2aCeRegWindow upper-bound check for the AllToAll CE registered path, preventing CE from being selected at message sizes beyond the arch-table ceiling.

ncclCeAvailable now has Float8 and ncclFloat8e5m2 checks for AllReduce

Test Changes

Tests Removed: UseCeAllReduce_MsgTooLargeReturnsFalse, DefaultOff_2Shot_Gfx1250_Isolated

UseCeAllReduce_MsgTooLargeReturnsFalse was a standalone test that called
rcclUseCeAr2Shot directly with a count exceeding the default 256 MiB cap and
expected it to return false. The test was removed because the message size upper
bound is no longer enforced inside rcclUseCeAr2Shot itself — it is checked
upstream in rcclSelectAllReduce via the twoShotWindow gate. Calling the helper
directly in isolation no longer exercises the right code path.

DefaultOff_2Shot_Gfx1250_Isolated was a parameterized case inside
RcclUseCeAllReduce_Isolated that asserted CE two-shot is off by default on
gfx1250 (because ceNonRegMax[AR]=0 in the gfx1250 table). The case was removed
because the fake rcclGetArchThresholds in the unit test environment always returns
nullptr, so the test was not actually exercising the gfx1250 table at all —
it was testing the null-table fallback path, which is covered by the existing
DisabledByDefault_Isolated case.

MessageTooLargeRejected_Isolated (inside RcclUseCeAllReduce_Isolated) was
similarly removed. It tested the size rejection via rcclUseCeAr2Shot with
baseEnv set, but the max-byte enforcement moved out of that helper into the
selector. The case would no longer have rejected the oversized message at the
intended point.

Tests Added: CE AllReduce Two-Shot Window Boundary Coverage

Issue Tracking

JIRA ID: AICOMRCCL-2422

Test Plan

CeAllReduceEligibilityTests (RcclCeNonRegMaxTab suite) — unit tests for the threshold query functions; updated to exercise both max and min bound overrides via the new env vars
wrap-test.cc host microtests — rcclSelectAllReduce integration path with forced env vars
CeInternalMPITests / CeMPITests — end-to-end CE AllReduce correctness under threshold overrides
Manual: all_reduce_perf -b 1K -e 8G -f 2 -g 8 with RCCL_CE_AR_2SHOT_MIN_BYTES /RCCL_CE_AR_2SHOT_MAX_BYTES set to verify window shifting takes effect and wrong=0

Test Result

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

therock-pr-bot · 2026-09-28T17:26:30Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/rccl/src/ce_coll.cc`, `projects/rccl/src/enqueue/enqueue.cc`, `projects/rccl/src/include/ce_coll.h`, `projects/rccl/src/include/rccl_arch_thresholds.h`, `projects/rccl/src/include/rccl_common.h` (+10 more); no test file found
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot

Copilot review overview

🟡 Changes recommended

Unresolved threshold, compatibility, selection, documentation, and LL128 gating issues remain.

Review effort: Lite
Findings: 4 · 2

Open (6)

LL128 AllReduce ignores RCCL_DDA_LL disable switch · New
Renamed environment variable breaks existing configurations · New
AlltoAll registered CE ignores CE_COLL_MAX_BYTES cap · New
CE 2-shot AllReduce ignores minimum-byte threshold · New
Staging default conflicts with documented memory settings · New
Diagnostic reports inverted symkRequested state · New

What changed in this PR

Adds runtime threshold controls for CE and symmetric kernels, increases CE staging capacity, and adjusts collective selection and DDA LL128 eligibility.

Changes:

Adds CE and symmetric-kernel threshold parameters.
Increases the CE staging default to 256 MiB.
Updates CE selection, diagnostics, and DDA LL128 eligibility.

File	Summary
`projects/rccl/src/rccl_wrap.cc`	Runtime thresholds, collective selection, and diagnostics
`projects/rccl/src/include/rccl_common.h`	New parameter declarations
`projects/rccl/src/include/ce_coll.h`	Increased CE staging capacity
`projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric_ll.cu`	LL128 eligibility changes

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

amd-argus · 2026-09-28T17:44:57Z

@cppccppccppc @sunguobao PTAL

amd-argus

ARGUS

Bottom line: Worth settling before approval: the Branch #3 || to && change at rccl_wrap.cc:1686 removes CE-registered AllGather entirely on gfx942/gfx950 and fails two in-tree host tests.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Findings

2 finding(s) that could not be posted as inline comments

projects/rccl/src/enqueue/enqueue.cc:4464 (Low) ceAgSymRegistered was not updated alongside AllGather branch #3, so the 'selection and dispatch agree' comment at enqueue.cc:4464 is now false The comment at enqueue.cc:4464-4466 says ceAgSymRegistered is the same predicate AllGather Branch #3 reports, but Branch #3 now also requires CTAPolicy & ZERO (rccl_wrap.cc:1685) while the twin at enqueue.cc:4467-4470 does not. The twin is live at enqueue.cc:4514 whenever decisionValid is false, which collectives.cc:159 and :189 both reach. Could the twin pick up the same term, or the comment be corrected?
projects/rccl/CHANGELOG.md:32 (Low) CHANGELOG.md:32 still names rcclCeAr2ShotMax(), the function this PR deletes CHANGELOG.md:32 still says the CE 2-shot cap is resolved at runtime via rcclCeAr2ShotMax(), which this PR deletes. It is the only remaining reference to that name in the tree. Could it be repointed at rcclCeNonRegMaxTab(table, ncclFuncAllReduce)?

Bottom line: Worth settling before approval: rcclCeNonRegMaxTab and extAlgoArchTable are declared in no header, so enqueue.cc:4447 does not compile.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

github-actions · 2026-10-05T21:43:17Z

No undocumented RCCL or NCCL environment-variable changes.

amd-argus

ARGUS

Bottom line: Worth settling before approval: five assertions in wrap-test.cc and CeAllReduceEligibilityTests.cpp still expect the gates this PR removed, including the test it rewrote at wrap-test.cc:5954.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Findings

1 finding(s) that could not be posted as inline comments

projects/rccl/test/host/wrap-test.cc:4758 (High) SelectAllGather_CeRegisteredViaSymmetricWindowsChosen at wrap-test.cc:4740 is broken by the branch #3 || to && change SelectAllGather_CeRegisteredViaSymmetricWindowsChosen at test/host/wrap-test.cc:4740 reached RCCL_CE_REGISTERED through the old || arm, on CTAPolicy=ZERO alone. Branch #3 now also requires rcclAllGatherCeRegisteredWindowTab, and the test passes nullptr buffers so g_symRegType stays ncclSymSendNonregRecvNonreg (test/host/fakes/sym_kernels_fakes.cc:21) and recvReg is false at rccl_wrap.cc:899. The assertion at test/host/wrap-test.cc:4758 fails. Could the test set g_symRegType = ncclSymSendNonregRecvReg?

Bottom line: Worth settling before approval: ceRegMax[AllGather] going to kThreshUnlimited on gfx942/gfx950 makes CE the default AllGather dispatch at any size through the uncapped ceAgSymRegistered arm.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Findings

1 finding(s) that could not be posted as inline comments

projects/rccl/test/host/wrap-test.cc:4758 (High) SelectAllGather_CeRegisteredViaSymmetricWindowsChosen cannot pass with Branch #3's new &&, and was not updated SelectAllGather_CeRegisteredViaSymmetricWindowsChosen at wrap-test.cc:4740 can no longer pass with Branch #3's new &&. The test passes nullptr buffers and never assigns g_symRegType, so the micro-test fake returns ncclSymSendNonregRecvNonreg (test/host/fakes/sym_kernels_fakes.cc:21) and rcclAllGatherCeRegisteredWindowTab returns false on the recvReg check at rccl_wrap.cc:899; previously the (CTAPolicy & NCCL_CTA_POLICY_ZERO) disjunct alone carried the assertion at wrap-test.cc:4758. The negative half at wrap-test.cc:4768 still passes, so the mutant guard survives. Could the test set g_symRegType = ncclSymSendNonregRecvReg before the call?

Bottom line: Worth settling before approval: SelectAllGather_CeRegisteredViaSymmetricWindowsChosen at wrap-test.cc:4758 can no longer pass with Branch #3's new &&, and was not updated with its sibling.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Findings

1 finding(s) that could not be posted as inline comments

projects/rccl/test/host/wrap-test.cc:4758 (High) SelectAllGather_CeRegisteredViaSymmetricWindowsChosen breaks on the branch #3 tightening and was not updated SelectAllGather_CeRegisteredViaSymmetricWindowsChosen (test/host/wrap-test.cc:4740) breaks on the branch #3 tightening. It never hooks g_getSymRegType, so the default fake yields ncclSymSendNonregRecvNonreg and rcclAllGatherCeRegisteredWindowTab returns false at rccl_wrap.cc:899, failing the EXPECT_EQ at test/host/wrap-test.cc:4758. The sibling test at :5938 got both compensations in this PR. Could this one get the same ScopedHook?

Bottom line: Worth settling before approval: the branch #3 tightening also breaks SelectAllGather_CeRegisteredViaSymmetricWindowsChosen at test/host/wrap-test.cc:4758, which this PR did not update.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Bottom line: Worth settling before approval: the AllGather Branch #3 && at rccl_wrap.cc:1742 breaks SelectAllGather_CeRegisteredViaSymmetricWindowsChosen at wrap-test.cc:4886, which precheckin runs.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Bottom line: Worth settling before approval: SelectAllGather_CeRegisteredViaSymmetricWindowsChosen cannot pass at this head, since the null-table fake makes rcclCeRegMaxTab return 0 and Branch #3's new && turns that into a miss.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Bottom line: Worth settling before approval: ddaLL128Max[AlltoAll] going to 0 at rccl_arch_thresholds.cc:49 removes the gfx1250 LL128 AlltoAll tier, and three existing tests assert the old value.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Bottom line: Nothing new worth raising; still open from earlier: ce_coll.h:28 and 1 more. Worth clearing before approval.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

Bottom line: Worth settling before approval: MinBoundRejects_Isolated expects false, but rcclUseCeAr2Shot reads no size bound at this head, so it returns true and the case fails.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

mustafabar

Do we really need start and end thresholds for env variables?
Perf debugging and tuning are 2 separate things. Tuning requires a start and end infra which you should have in your arch table. Env variables are for debugging perf.
So why not just make RCCL_FORCE_CE and RCCL_FORCE_2SHOT_CE work for us to understand their perf. At the same time, surface the selection correctly at rccl-tests and log why it wasn't forced for whatever reason. The staging buffer size requirement and environment are OK.

isaki001 · 2026-10-12T14:09:34Z

Do we really need start and end thresholds for env variables? Perf debugging and tuning are 2 separate things. Tuning requires a start and end infra which you should have in your arch table. Env variables are for debugging perf. So why not just make RCCL_FORCE_CE and RCCL_FORCE_2SHOT_CE work for us to understand their perf. At the same time, surface the selection correctly at rccl-tests and log why it wasn't forced for whatever reason. The staging buffer size requirement and environment are OK.

I agree with needing env. vars for debugging instead of tuning.
PR https://github.com/ROCm/rocm-systems/pull/13098/changes addresses this for other algos.
I didn't do it for CE, because that is what -x 2 will do, force CE everywhere so an env. var for force-CE seemed redundant.
RCCL_FORCE_CE has a different use currently for enabling CE when symmetric buffers aren't there, so I didn't want to change how it is used.
I see your point with the min/max env. vars seeming like an overkill while an force-env would be sufficient.
However, when I did the tuning infra, I wanted to allow this flexibility to give a way out of current tuning if needed without having to force a particular algo everywhere.
We shouldn't need these env. vars outside of debugging, but they seemed like a nice fail-safe to have without much overhead in the code.
They may also help with unit-testing and eval of other tuning tools.
They fit well with current infra, since they just override the tuning-table threshold in one routine, so we can easily remove them as well.

…f cta policy and force flag

amd-argus

ARGUS

This PR changes 1,055 lines across 22 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: LGTM.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

This PR changes 1,059 lines across 22 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: LGTM.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

…ts (#13819) JIRA ID: AICOMRCCL-2826 ## What Fixes the host-test failures on develop: 5 tests in `rccl-UnitTestsMicroDdaColl` and 1 in `rccl-UnitTestsMicroRocmwrap`. ## How #12386 removed the `RCCL_DDA_LL` check from the fabric LL128 eligibility functions (LL128 is gated by `RCCL_DDA_LL128` upstream). #13585, merged a few hours earlier, still expected `RCCL_DDA_LL=0` to reject LL128. The tests now assert `RCCL_DDA_LL=0` leaves LL128 eligible. Develop also reworded the cuMem driver-version warning ("requires ROCm >= 7.12 or the 7.0.2.x backport"); `RocmwrapMicrotest` now looks for that text. ## Test Plan `rccl-UnitTestsMicroDdaColl` 271/271 and `rccl-UnitTestsMicroRocmwrap` 35/35, each 20x shuffled clean (standalone build). Re-adding the old `RCCL_DDA_LL` gate fails the new test. --------- Co-authored-by: Claude <noreply@anthropic.com>

systems-assistant · 2026-10-21T13:31:58Z

TheRock Submodule Bump Activity

Newest first

2026-10-21 — Included in TheRock via ROCm/TheRock#9398.

## Motivation `ncclCommRevoke` and `ncclCommAbort` set `abortFlag` so in-flight kernels can drain. Generic device primitives already sample that flag, but the standalone DDA kernels did not. On gfx1250, an incomplete DDA collective (for example, a rank that never launches its AllReduce) left the other ranks spinning on peer flags or barrier mailboxes that would never arrive. Revoke and abort then hung in stream sync (`RevokeMPITest.IncompleteCollective_Revoke_Shrink_Collective`). DDA IPC on gfx942/gfx950 has the same shape in `IpcGpuBarrier`. ## Technical Details Every DDA wait now polls `comm->abortFlagDev`, at the same cadence as `checkAbort` (`kDdaSpinsBeforeCheckAbort`, in the new `dda_kernel_params.h`). This covers the LL and LL128 flag waits in AllReduce (one-shot and two-shot), AllGather, ReduceScatter and AlltoAll, and the fabric and IPC GPU barriers. Launchers pass a `dda::common::DdaKernelParams` by value to the kernels, wait helpers and barriers. It holds the abort flag today, and later per-launch values go there rather than as extra kernel arguments. On abort a wait gives up rather than returning early from the kernel where that would skip shared work: polls `break` so the LL epoch update and `__syncthreads` still run. The LL two-shot AllReduce uses a sticky abort, skipping the remaining waits once the flag is seen, because an early return after each unrolled wait nearly doubled the 8-rank kernel and slowed the single-node two-shot path. A thread that has aborted does not publish its phase-2 sums, which are partial, so a peer that has not seen the abort cannot accept them as the reduced shard. `IpcGpuBarrier` and `FabricGpuBarrier` keep a sticky `aborted_` for the rest of the kernel. Once a signaling lane gives up a put or a wait, the block signals no later barrier, so a late peer cannot pair its earlier flag with this block's next one and read a two-shot shard this rank never finished. The signaling lanes all sit in wave 0, so they share the result through `__any`; `FabricGpuBarrier` uses `__syncthreads_or` when the clique is wider than one wave. This replaces a load of the pinned-host abort flag before every put. DDA stays enabled when revoke is used, because revoke has to stop kernels that are already running. A collective issued on a revoked comm is rejected: DDA, CE 2-shot and GIN-SDMA return from the collective before `ncclEnqueueCheck`, so `rcclCollectiveMustUseEnqueuePath` now sends a revoked comm through enqueue, which returns `ncclInvalidUsage` as `nccl.h` documents. Without this, a collective on the revoked parent would run DDA again on mailboxes left with unconsumed flags. A parent comm shrunk with `NCCL_SHRINK_ABORT` stops using DDA (`comm->ddaDisabledAfterAbort`, checked in `rcclDdaEnabled`). The abort can leave ranks with different LL epochs and unconsumed barrier mailbox slots, so a later DDA collective on that parent could hang or return a peer's stale data. The shrunk child does not share the parent's resources and is not affected. Tests (the MPI cases are in `RevokeMPITests.cpp`): - `RevokeDdaMPITest.IncompleteCollective_*` (11 cases, gfx1250): one rank skips a DDA collective, the rest revoke and shrink. They cover LL and LL128 one-shot and two-shot AllReduce, the VMM barrier, and LL/LL128 AllGather, ReduceScatter and AlltoAll. Each case first confirms through `rcclGetCollImplInfo` that the intended DDA path is selected, and after the revoke checks that the same collective on the revoked parent returns `ncclInvalidUsage`. - `RevokeDdaIpcMPITest.IncompleteCollective_DdaIpc*` (5 cases, gfx942/gfx950, 8 ranks): the same pattern for DDA IPC AllReduce (flat and the three-barrier tree kernel), AllGather, ReduceScatter and AlltoAll. - `AbortDdaMPITest.IncompleteCollective_*` (5 cases): the same pattern with `ncclCommAbort`, for LL one-shot and two-shot AllReduce, and AllGather, ReduceScatter and AlltoAll LL. - `ShrinkAbort_DisablesDdaOnParent` in both `RevokeDdaMPITest` and `RevokeDdaIpcMPITest`: after `NCCL_SHRINK_ABORT`, the parent no longer selects DDA and the child's AllReduce succeeds. - `DdaIpcGpuBarrierOrdering.AbortedBarrierStaysSilent` (one GPU): a barrier that gave up on abort must not signal the next one after late peers drain the first. It fails with the sticky `aborted_` removed. - `Rcclwrap.EnqueueGuards_RevokedForcesEnqueue`: a revoked comm takes the enqueue path. - The DDA fabric launcher microtests check that `comm->abortFlagDev` reaches the kernel; the mock comm now carries a non-null flag so the check can fail. Launch, abort and stream-drain results are voted across ranks, and post-revoke waits are bounded, so a failure on one rank fails every rank instead of hanging. The test runner configs for MI300, MI355X and MI455 include the new cases. ### Known gaps These are out of scope for this PR and tracked as follow-ups: - **`ncclCommRevoke` and `NCCL_SHRINK_ABORT` with a stuck DDA kernel.** Both set `abortFlag`, sync only `hostStream` and `deviceStream`, then clear the flag (`commRevokeAsync` and `ncclCommShrink`). DDA kernels run on the user stream, which `deviceStream` is not ordered after outside graph capture, so a stuck kernel that does not poll the flag inside that window keeps spinning. In practice the window is long enough for the poll (every `kDdaSpinsBeforeCheckAbort` spins), which is why the revoke tests pass, but it is not guaranteed. Native single-stream kernels and NCCL upstream have the same window. `ncclCommAbort` leaves the flag set and is not affected. - **DDA stays off on a parent shrunk with `NCCL_SHRINK_ABORT`.** Turning it back on needs a cross-rank resync of the LL epochs and barrier mailboxes. - **Graphs captured before the abort.** The DDA gate applies at enqueue, so a graph that captured a DDA kernel before the abort still replays it. - **Kernels that start after the flag is cleared.** A DDA kernel that is queued behind other user-stream work and only starts after `ncclCommRevoke` or `NCCL_SHRINK_ABORT` clears `abortFlag` never sees the abort, and hangs if a peer is missing. ## Issue Tracking JIRA ID: ROCM-33096 (gfx1250, DDA fabric), ROCM-32381 (gfx942, DDA IPC). ## Test Plan Run each `RevokeDdaMPITest` case in its own process, as the runner configs do. In one process the first case caches `RCCL_DDA_LL` and the LL sub-thresholds, so later `MpiEnvGuard`s have no effect and some cases skip. `AbortDdaMPITest.*` sets no overrides and can run in one process. - `mpirun -np 4 ./rccl-UnitTestsMPI --gtest_filter=<case>` for each `RevokeDdaMPITest`, `RevokeMPITest` and `RevokeNonBlockingMPITest` case in the `revoke_mpi_tests_4rank` group (for example `--gtest_filter=RevokeDdaMPITest.IncompleteCollective_DdaVmmBarrier`), and `--gtest_filter='AbortDdaMPITest.*'`. - The same MPI cases at 8 ranks across two nodes (`mpirun -np 8 --host <nodeA>:4,<nodeB>:4 ...`). - `mpirun -np 8 ./rccl-UnitTestsMPI --gtest_filter=RevokeDdaIpcMPITest.<case>`, one per process, on 8 gfx942/gfx950 GPUs, plus the revoke, shrink, abort and grow suites of `mi300x_mellanox_ib.json`. - `./rccl-UnitTestsFixtures --gtest_filter='DdaIpcGpuBarrierOrdering.*'` and `./rccl-UnitTestsFixturesDebug --gtest_filter='Rcclwrap.EnqueueGuards_*'`. - Host microtests: `rccl-UnitTestsMicroDdaColl` and `rccl-UnitTestsMicroDda`. - rccl-tests with data checking (`-c 1`) for AllReduce, AllGather, ReduceScatter and AlltoAll on one and two nodes, plus two-node AlltoAll and ReduceScatter with `RCCL_DDA_LL128_THRESHOLD=67108864`. `NCCL_DEBUG_SUBSYS=COLL` logs confirm the DDA path at each size. - rccl-tests performance against the develop base on one and two nodes: the four collectives from 1 KB to 64 MB (32 MB on two nodes), plus AllReduce forced onto the LL two-shot kernel (`RCCL_DDA_LL_ONESHOT_THRESHOLD=0`). ## Test Result **gfx942, DDA IPC** (8x MI300X, one node, ROCm 7.13.0 rc2, DebugFast, `7bb78ac42da`): - The revoke, shrink, abort and grow suites of `mi300x_mellanox_ib.json`: 62 pass, 2 fail and 19 skip on this hardware, including every `RevokeDdaIpcMPITest.IncompleteCollective_*` case and `DdaIpcGpuBarrierOrdering.*`. `RevokeDdaIpcMPITest.ShrinkAbort_DisablesDdaOnParent`, added after that run, passes on its own. The 2 failures are `SymAbort_Kernel.LLAllToAll_AbortReleasesKernel` and `SymAbort_Kernel.LsaBarrier_AbortReleasesKernel`, which time out in communicator setup on develop as well. - `Rcclwrap.EnqueueGuards_*` passes. With the sticky barrier abort or the revoked-comm divert reverted, `AbortedBarrierStaysSilent`, `EnqueueGuards_RevokedForcesEnqueue` and the IPC AllReduce revoke cases fail. - Host microtests: `rccl-UnitTestsMicroDda` passes 51 of 51. `rccl-UnitTestsMicroDdaColl` passes 266; 5 LL128 eligibility tests fail on develop as well, after #12386 removed the `RCCL_DDA_LL` gate from LL128 eligibility. - With the barrier fix applied on `cc2700e596d`, rccl-tests `-c 1` reports 0 wrong values for all four collectives on DDA IPC, and performance against `cc2700e596d` is within noise from 1 KB to 8 MB (3 reps). **gfx1250, DDA fabric** (4 GPUs per node, `ctheliosp-rck-g02-j11-12` alone and with `j11-11`, one MNNVL fabric clique): - At `65410fcff30` (the barrier fix, before the revoked-comm divert and the new tests), 80 of 86 cases in the revoke, shrink and abort suites pass on one and two nodes. Of the 6 failures, 3 two-node cases did not start (`ORTE was unable to reliably start one or more daemons`), 2 are the `SymAbort_Kernel` timeouts above, and 1 is a stale `SymCheckMode_Registration` filter name in the mi455 config on develop. rccl-tests `-c 1` reports 0 wrong values in every run that completed. - `7bb78ac42da` is not yet tested on gfx1250: the nodes moved to a new ROCm build mid-run and are now in use. Results to follow. - Earlier measurements against develop `dfb3f8c9e28` (median latency, 2% threshold with non-overlapping ranges): - At `134defe6cc6`, AllReduce forced onto the LL two-shot kernel is 3–6% slower (about 1–1.4 µs) from 64 KB to 2 MB, on one and two nodes. At default thresholds this kernel is used at 2 MB only. - Two-node ReduceScatter on DDA fabric LL is 2–3.4% slower from 32 KB to 2 MB; on DDA fabric LL128 it is 5–7% faster at 4–8 MB and about 4% faster at 16–32 MB. - Everything else is within noise, apart from single-point 3% differences: one-node AllReduce 2 KB slower, two-node AllGather 32 KB faster and AlltoAll 2 KB slower. - With the two-shot review fixes (`8fecb7a64bb`), the LL two-shot AllReduce is within -1.5% to +0.2% of `134defe6cc6` at every size from 64 KB to 2 MB, on one node (4 ranks, 5 reps) and two nodes (8 ranks, 3 reps). ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Nikita Gusev <nikita.gusev2@amd.com>

…the RCCL 2.32.3 for ROCm 10.3.0 (Unreleased) section (#14929) This PR does two things to `projects/rccl/CHANGELOG.md`. 1. Removes the duplicate `## RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` heading. #13544 removed it, and #13128 brought it back. The section is not a pure copy: 48 of its entries have no identical copy elsewhere. 47 of them are folded into `## RCCL 2.32.3 for ROCm 10.2.0` instead of being deleted, and the 48th is the older wording of the `RCCL_CE_AR_STAGING_BYTES` entry, which is dropped in favor of the newer one. Where the same entry appeared in both sections, one copy is kept. 2. Moves the entries below, which sit in release sections but whose code was merged to develop after the `rocm-systems` commit that TheRock `release/therock-10.2` pins (`ded49fd`) and is not in that commit. This moves 53 of them into `## RCCL 2.32.3 for ROCm 10.3.0 (Unreleased)`, under the matching subsections. Not moved: 10 entries that were attributed to #13544 (the CHANGELOG-only PR that last touched them), so the code PR could not be traced. They stay in the 10.2.0 section and are unverified, so the RCCL owners should confirm them. They begin: `RCCL_TDM_LL128_ENABLE`, `Ll128TdmStoreTest`, the TDM reduction path for SIMPLE, Anvil SDMA strong GIN signals, VA GIN signals on Anvil SDMA, `NCCL_P2P_DISABLE=1` on shared topology nodes, queue evictions at process exit, multi-receive `irecv` with `NCCL_IB_USE_INLINE=1`, Anvil SDMA signal cells across logical contexts, and GIN signal shadow sizing. - #13328 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_BCM_LINKS_PATH`: fallback directory path for Broadcom PEX890xx multi-host switch inter-switch-link topology. On systems where a single physical PEX switch is partitioned across multiple CPU domains (upstream ports), RCCL needs to know which ports share the same switch so it can classify cross-partition PCIe paths as `PXB` instead of `PHB` and enable GDR across them. The `switch_discovery` kernel module exposes this via sysfs (`/sys/kernel/pci_switch_link/virtual_switch_links`); when that path is not available, RCCL reads the same directory layout from `RCCL_BCM_LINKS_PATH`. The companion userspace tool `rccl_bcm_isl_discover.sh` discovers the switch groups by reading the PCIe Device Serial Number (DSN) from extended config space (vendor `1000`, device `c030`) and populates this directory. - #12742 (merged 2026-10-18, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_MLOPART`: allows GIN on MLO/CPX partitions. A partition reaches xGMI and the network over the physical device's paths, so GIN is as available to a partition as to the whole GPU. By default a single-node partitioned communicator keeps GIN, and one spanning several nodes does not, because past one node admitting a partition also reaches the hierarchical Copy Engine path. Set `1` to allow GIN on partitions on any number of nodes, or `0` to restore the previous partition exclusion. - #3894 (merged 2026-10-16, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_TREE_COMPACT` (default `0`): builds the inter-node trees with a half-interleave node remapping, for jobs whose nodes are ordered by network domain. Both trees keep cross-domain hops at the leaves only when the nodes split evenly into two domains; with an uneven split, the first tree does so when the first domain holds at least half of the nodes, and the second tree still crosses domains. All ranks in a communicator must use the same value; initialization fails with `ncclInvalidUsage` otherwise. - #12856 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_FABRIC_TELEMETRY_ENABLE` (default `0`): samples every amd-smi fabric telemetry category for the GPUs this process opens communicators on and logs the counters that moved, plus a closing report over the whole sampled span when a device's last communicator is destroyed. Diagnostic only, and off by default. Also needs `RCCL_USE_AMD_SMI_LIB=1` and `render`-group access to the IFoE config character device; without either, RCCL logs why telemetry was skipped and initialization continues. - #12856 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_FABRIC_TELEMETRY_INTERVAL_MS` (default `1000`): sampling period for `RCCL_FABRIC_TELEMETRY_ENABLE`, clamped to a 100 ms floor and a one-hour ceiling. A non-positive value disables telemetry entirely. - #12573 (merged 2026-10-14, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Hierarchical symmetric `ncclAllGather` for multi-node using GIN and LSA (`AllGather_RailRing_LsaST`): a GIN ring over the rails team moves data between nodes and LSA stores broadcast it to the other GPUs in each node. - #14140 (merged 2026-11-01, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Hierarchical symmetric `ncclAllGather` for partition mode in gfx1250 scale-up domain using LSA (`AllGater_HierLsaST`): each rank stores its data once to its counterpart partition on every other GPU, which forwards it to the other partitions there, so data crosses to each GPU once instead of once per partition. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_ALLREDUCE_GFX1250_ENABLE` (default `0`): opt in to GinScatter AllReduce on gfx1250 for symmetric sum messages from 64 MiB through 4 GiB. Left off, gfx1250 keeps the pre-existing AllReduce selectors. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_ALLREDUCE_SCRATCH_BYTES` (default 512 MiB): GinScatter incoming scratch. `0` skips the allocation and GinScatter is not selected. The buffer is allocated only on the GinScatter path. - #13416 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table is currently a copy of the SPX table, so tuning is unchanged. - #13727 (merged 2026-11-03, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_UDMA_POLICY` (default `0`): selects how QPs are assigned to the two uDMA pipelines of an AINIC with the `IB-CAST` transport. `0` pins each channel's send QPs, and separately its receive QPs, to one pipeline per device, `1` round-robins P2P QPs, `2` leaves the choice to the driver, and `3` splits each connection across both pipelines (the default of two QPs per connection is enough; splitting is lost with one QP per connection and on RMA comms). Out-of-range values log a warning and fall back to `0`. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_AR_2SHOT_MIN_BYTES` (default `-1`): overrides the minimum message size for the CE 2-shot AllReduce path. When `-1`, resolved from `ceNonRegMin[AllReduce]` in the per-arch table. - #12310 (merged 2026-10-28, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_REDUCESCATTER` (default `0`): opt-in Copy Engine ReduceScatter on a single-node symmetric communicator with `CTA_POLICY_ZERO`. Always stages through `ceARTmpBuf`; user-buffer registration is not required. No message-size cap (large messages are pipelined). - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_COLL_MAX_BYTES` (default `-1`): overrides the CE registered-window and 2-shot size cap for all non-AllReduce collectives (AllGather, AlltoAll, etc.). When `-1`, resolved from `ceRegMax[func]` or `ceNonRegMax[func]` in the per-arch table. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_COLL_MIN_BYTES` (default `-1`): overrides the minimum message size for CE non-AllReduce paths (both registered-window and scratch). When `-1`, resolved from `ceNonRegMin[func]` in the per-arch table. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_SYM_K_MAX_BYTES` (default `-1`): overrides the symmetric-kernel upper bound for all collectives. When `-1`, resolved from `symMaxR2[func]` (eager) or `symMaxR2Graph[func]` (graph-capture) in the per-arch table. `kThreshUnlimited` (SIZE_MAX) in the table means no suppression. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_SYM_K_MIN_BYTES` (default `-1`): overrides the symmetric-kernel lower bound for all collectives. When `-1`, resolved from `symMinR2[func]` in the per-arch table. Below this size, DDA is preferred over symk. - #13995 (merged 2026-10-27, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_GDR_FLUSH_FLAG` (default `0`): replaces the GPUDirect RDMA flush of Simple-protocol receives in the `IB` network transport. Instead of the proxy fencing each receive with an RDMA read and waiting for it, the sender writes a per-QP flag into GPU memory after the payload without relaxed ordering, and the GPU waits for it. A communicator uses it only when every rank enables it and can use it (gfx942 or gfx950, the `IB` transport, and no adaptive routing, out-of-order receive queues, ID-based receive matching (`NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=1`) or port-failover resiliency (`NCCL_IB_RESILIENCY_PORT_FAILOVER`)); otherwise every rank keeps the RDMA read flush. Connections over a fused NIC (NIC Fusion, including a fused vNIC chosen by subnet-aware routing), with more than 8 QPs per request, or that cannot register the flag memory also keep it. On gfx950 with 2 to 8 nodes of 8 GPUs, such a communicator also selects Ring/LL128 from 16 MB per node, Tree/Simple from 32 MB per node and Ring/Simple from 64 MB per node for `ncclAllReduce`, and switches `ncclReduceScatter` from LL128 to Simple above 6 MB per rank instead of 8 MB, unless `NCCL_ALGO` or `NCCL_PROTO` is set. On MI355X it raises large `ncclAllReduce` and `ncclReduceScatter` bandwidth by 10 to 22% and 1 to 8 MB `ncclAlltoAll` by 4 to 12%. - #13995 (merged 2026-10-27, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_GDR_FLUSH_FLAG_TUNE` (default `0`): additional tuning for communicators that use `RCCL_IB_GDR_FLUSH_FLAG` on gfx950 with 8 GPUs per node on 2 to 8 or 16 nodes. Ring and tree connections get a 4x LL128 buffer (send/recv keeps the default, and `NCCL_LL128_BUFFSIZE` disables it); `ncclAllReduce` selects Ring/LL128 from 12 MB per node up to what one ring loop carries, then Tree/Simple up to 64 MB per node and Ring/Simple on 48 channels up to 128 MB per node, and Tree/Simple from 8 to 64 MB per node on 16 nodes; `ncclReduceScatter` keeps LL128 up to 7 MB per rank, then Simple, on 16 nodes too. All ranks must use the same value. - #13757 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_PCIE_ORDERED_RING` (default `1`, also accepted as `NCCL_PCIE_ORDERED_RING`): pins the ring graph search to the order GPUs appear in the topology array, so every channel uses the same ring. Without it the search picks a different permutation per channel, so a device exchanges with many different peers; pinning the order reduces fan-out to one peer per direction per device across all rings, which cuts the connections, FIFO buffers and setup work each device carries. On gfx1250 partitioned (DPX) GPUs the unpinned search also mixes scale-up and scale-in hops between channels; pinning keeps the two partitions of a die adjacent, though ring throughput is bounded by its slowest hop either way. This constrains GPU visitation order only: path validation, bandwidth constraints and backtracking are unchanged. It applies to gfx1250 ring graphs on a single host where every rank is a GPU of that topology, so trees, CollNet, other architectures and multi-host communicators are unaffected, and `NCCL_GRAPH_FILE` and `NCCL_RINGS` still take precedence. Set `0` to restore the unconstrained search, or `2` to fail communicator initialization instead of falling back when no ordered ring passes validation. - #10885 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_MULTI_SEGMENT_REGISTER` default changed `1` -> `0` on HIP and stays `1` on other builds. On HIP, multi-segment registration stays off until a job sets the variable. Set `1` to register one DMA-BUF memory region per physical segment where the transport supports it. - #10885 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * NET registration counts physical segments over the whole registered buffer, not only the range a collective touches. With `NCCL_MULTI_SEGMENT_REGISTER=0`, a collective on any part of a multi-segment registration skips NET registration and stages, instead of retrying a single-region registration on every collective. - #10891 (merged 2026-10-25, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * GIN window registration on HIP fails with `ncclInvalidUsage` when the window spans several device-memory segments, instead of registering a memory region that covers only the first segment, and when it contains host VMM memory and `NCCL_DMABUF_ENABLE=0`. - #12947 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `install.sh --all_unrolls` (`-DBUILD_ALL_UNROLLS=ON`) now adds the unroll factors the multi-arch default skips (8 and 16) on top of each architecture's existing set. It no longer replaces that set or clears the gfx1250 pin on unroll 32. The flag stays off by default, so a normal build does not compile unroll 8 or 16. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The GIO topology model is applied only to systems with exactly 16 GPUs. It was previously applied to any system with 16 or more GPUs, where the 16-GPU model does not describe the topology; such systems now fall through to the other models. - #13581 (merged 2026-10-19, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Enabled the single-node one-slice optimization for gfx1250, which was previously limited to gfx942 and gfx950. Ring Simple AllGather, AllReduce and ReduceScatter now use one slice per chunk instead of two, halving the per-chunk peer handshakes. The gain is largest below about 128 MB and tapers to a few percent at multi-GB sizes, where the ring is already near the link bandwidth. LL, LL128 and non-ring algorithms are unaffected. - #14123 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * On gfx1250, the DDA LL128 `ncclAllGather`, `ncclAllReduce` and `ncclReduceScatter` receivers keep each line's payload from the cooperative load that saw its flag instead of reading the slice a second time. The DDA scratch bank is now floored to 128 bytes instead of 16, so an LL128 wire line never straddles two 128-byte lines when `RCCL_DDA_FABRIC_BUFFER_SIZE` sets an unaligned size. Results are unchanged. Other architectures are unaffected. - #13415 (merged 2026-10-19, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `fp8_e4m3` and `fp8_e5m2` sum reductions on gfx1250 convert eight elements per instruction using the packed conversion instructions, instead of two. Results are unchanged. Other architectures and reduction operators are unaffected. - #13406 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The rocSHMEM GDA GIN backend (`NCCL_GIN_TYPE=6`) now sets the RoCE traffic class of its queue pairs from `NCCL_GIN_IB_TC`, then `NCCL_IB_TC`, then the communicator traffic class, matching the IB proxy backend. It previously read `ROCSHMEM_GDA_TRAFFIC_CLASS`, which no longer affects these queue pairs. The default remains traffic class 0 when none of these is set. - #13357 (merged 2026-10-23, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_SYM_TMA_ENABLE` now defaults to `1`, matching NCCL 2.31.2, so on gfx1250 (MI450) the Tensor Data Mover symmetric kernels (`AllGather_TmaST`, `AllReduce_RSxTmaLD_AGxTmaST`, `ReduceScatter_TmaLD`) are offered to the symmetric cost model whenever an application uses symmetric registered memory. The cost model selects them only where their deep loop engages, and launches them at the widest allowed block count. They require 16-byte aligned buffers and a build with gfx1250 in `GPU_TARGETS` and a ROCm SDK that ships the TDM descriptor header. For in-place AllGather and ReduceScatter the per-rank slice must also be a multiple of 16 bytes, so that every rank selects the same kernel. Other architectures are unaffected. Set `NCCL_SYM_TMA_ENABLE=0` to restore the previous behavior. - #13357 (merged 2026-10-23, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * rccl-tests `-A 1` (and `rcclGetCollImplInfo` / `rcclSymKGetInfo`) now report a symmetric collective that selects a Tensor Data Mover kernel as `SymTDM` instead of `SYM`. The symmetric reporting query also accounts for buffer alignment, so it no longer omits the TDM kernels that the live dispatch would pick. - #12291 (merged 2026-10-21, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The AMD Pensando 'ionic' RDMA provider for the rocSHMEM GDA GIN backend (`NCCL_GIN_TYPE=6`) now properly enforces the required support for creation of 'collapsed' single-entry completion queues. If network driver and firmware support for this feature is unavailable, an initialization failure will occur rather than a silent fallback to plain completion queues. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_DDA_LL128_RS_THREADS` removed (was default `1024`): reduce-scatter LL128 now launches a fixed 512-thread block, one warp per slice, so the thread count is no longer a tuning knob. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, causing an `ncclGather` performance regression on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed LL-protocol latency on gfx950 when RCCL is built against a ROCm older than 7.0.2 without the 128-bit global load/store builtins (for example ROCm 7.0.1, or `-DDWORDX4_INTRINSICS=OFF`). Every store to the user output buffer was followed by a release fence and an L2 writeback, which made 4 KB to 64 KB `ncclBroadcast` about 30% slower on MI355X. Builds against ROCm 7.0.2 or later were not affected and are unchanged. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed a 15 to 30% `ncclScatter` bandwidth loss from 16 MB upward on multi-node AINIC jobs, caused by the move to two network channels per peer (for example on 4 nodes of 8 GPUs). Scatter now uses and connects the most channels per peer whose flows still fit the P2P channel pool. Point-to-point operations and the other collectives keep their channel count. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed `ncclCommAbort` and `ncclCommRevoke` hanging when a DDA collective (IPC on gfx942/gfx950, fabric LL/LL128/VMM on gfx1250) was waiting on a peer that never joined. DDA kernels now observe the communicator abort flag while waiting. A parent communicator shrunk with `NCCL_SHRINK_ABORT` no longer uses DDA for later collectives. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * DDA LL128 eligibility functions (`ddaLL128ArOneShotEligible`, `ddaLL128ArTwoShotEligible`) no longer check `RCCL_DDA_LL`. LL and LL128 are now independently gated: `RCCL_DDA_LL=0` disables only the LL tier; `RCCL_DDA_LL128=0` disables only the LL128 tier. Both checks are enforced via `rcclDdaLLThresholdTab` / `rcclDdaLL128ThresholdTab` at the selector level, which return 0 when the respective param is disabled. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * DDA LL128 is disabled for ReduceScatter and AlltoAll on gfx1250: `ddaLL128Max[RS]` and `ddaLL128Max[A2A]` are 0 in the gfx1250 arch table. - #13005 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * With `NCCL_CTA_POLICY=2` (`NCCL_CTA_POLICY_ZERO`), AllReduce on registered symmetric windows now takes the registered CE path whenever it is available, even when the symmetric kernel is eligible, matching AllGather and upstream NCCL. A non-zero arch-table `ceRegMax[AllReduce]` no longer caps it under this policy; a zero entry and an explicit `RCCL_CE_AR_REG_MAX_MSG_BYTES` still apply. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_CE` (default `1`): acts as a prerequisite gate for CE-scratch selection. When `RCCL_FORCE_CE=0`, CE-scratch is disabled entirely for all collectives regardless of other flags. When non-zero, CE-scratch fires if `CTA_POLICY_ZERO` (`-x 2`) is set (bypassing the size window) or if the message size falls within the arch-table `ceNonRegMax` window. AlltoAll CE-scratch now has a size window (`ceNonRegMax[AlltoAll]` in the per-arch table, currently 0 on all arches). - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Registered-window CE AllReduce is now eligible without requiring ``NCCL_CTA_POLICY=2`` > (``CTA_POLICY_ZERO``); it is gated only by the per-arch table window and ``RCCL_CE_ALLREDUCE``. > ``RCCL_FORCE_CE_ALLREDUCE`` now applies exclusively to the non-registered (2-shot/staged) CE > AllReduce path. - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_SYMMETRIC` (default `0`): when set to a non-zero value, biases the `ncclAllReduce` > backend selector to prefer the symmetric kernel path over all other candidates. The selector > still requires the operand buffers to be registered symmetric windows; if the symmetric > path is ineligible the selector falls back normally. - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_DDA` (default `0`): when set to a non-zero value, biases the `ncclAllReduce` > backend selector to prefer DDA (Direct Data Access) paths over all other candidates, > including symmetric and Copy Engine. If no DDA path is eligible for the given message > size and communicator the selector falls back normally. - #14254 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed an intermittent hang, and possible wrong results, in the first DDA fabric LL all-reduces after communicator initialization on gfx1250 (`NCCL_MNNVL_ENABLE=1`). Initialization zeroed the LL scratch buffer and epoch counters with an asynchronous null-stream memset that could land after a rank's first LL kernel, resetting its epoch so that rank used a different flag and scratch bank than its peers. Initialization now waits for the zero fills, and the ranks agree that every rank zeroed its buffers before any rank can launch an LL kernel. If any rank fails to set up its LL buffers or GPU barrier, every rank disables the fabric DDA path together. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed memory corruption, uninitialized reads and a division by zero when parsing malformed or oversized `NCCL_RINGS` values. A ring list with more channels than `MAXCHANNELS`, a ring naming more GPUs or more NET endpoints than supported, or a GPU or NET number too large for an `int` wrote past the channel and parser arrays or wrapped around; NET slots a ring did not name were copied into the graph uninitialized; and input from which no channel could be parsed divided the system bandwidth by zero channels. Such input is now rejected with `ncclInvalidUsage`, and bandwidth is zero when no channel is parsed. Parsing a `NULL` ring string returns `ncclInvalidArgument` instead of dereferencing it. A ring that names its NETs between its GPUs rather than at its ends, for example `"0 N0 1 2 N1 3"`, no longer stores a negative NET id in the graph, which failed multi-node initialization; such a channel is ignored. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed a crash and silent misconfiguration when parsing malformed `RCCL_TREES` values. A tree not terminated by `)`, for example `"(1)(2)(3)(4)(5)(6"`, searched past the end of the string and segfaulted; a GPU number larger than an `int` wrapped around and built a tree over an unintended GPU, as `"(4294967299)"` did over GPU 3; and a stray non-digit character resolved to a GPU number left over from the previous tree, so `"(5)(6)|(0)(x)"` parsed as `"(0)(6)"`. The number of trees, the GPUs per tree, the transcribed tree length and the GPU index used for remapping are now bounded. Invalid input is rejected with `ncclInvalidUsage`, and a `NULL` tree string returns `ncclInvalidArgument`. - #13815 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed wrong results and hangs in application kernels that use the GIN device API, such as the `07_kernel_fusion` GIN and hybrid RMSNorm examples. `nccl_device/coop.h` selected `__syncwarp` from `ROCM_VERSION`, which only RCCL's own build defines, so in application code the thread and warp cooperative groups synchronized the whole block with `__syncthreads()`. The GIN barrier calls them from a subset of threads, which left the block's waves out of step. The headers now fall back to the HIP version when `ROCM_VERSION` is not defined. - #12743 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed the GIN plugin never being finalized for a communicator that had connected a GIN backend, leaking the plugin's per-communicator context and leaving its reference count one too high after every `ncclCommDestroy`. `ncclGinHostFinalize` closed the GIN connections and then zeroed the whole `ncclGinState`, so the `ncclGinFinalize` that follows it in `commFree` saw no active backend and skipped the plugin `finalize()` call. Host finalize now clears only the per-connection state and hands the backend records to `ncclGinFinalize`, which finalizes each one on the last shared-resource reference. `commFree` returns the first of a `closeColl` or plugin-finalize error after those frees; if both fail, the `closeColl` error is the one reported. An `ncclRmaProxyFinalize` error still aborts the rest of teardown immediately. On the Anvil SDMA backend (`NCCL_GIN_TYPE=7`) this was one leaked init context per connected communicator. - #12679 (merged 2026-10-16, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed a crash (`HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`) on gfx110x and gfx120x when the Tree algorithm runs over shared memory with more channels than the Tree graph search found, for example with `NCCL_ALGO=Tree` and `NCCL_P2P_DISABLE=1`. Channels added after `ncclTopoPreset` had no intra-node tree links; `ncclTopoPostset` now sets the tree and CollNet chain links on every channel, as it already did for rings. - #11133 (merged 2026-10-14, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed the host one-sided RMA API being unavailable on multi-node communicators whose ranks had fused NICs. NIC fusion merges devices for the regular transport only; GIN and RMA select from their own topology node types and never receive a fused device. - #12750 (merged 2026-10-18, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed Anvil SDMA GIN small-message AllToAll latency after `ncclGinFenceLevel` support. A sub-threshold `gin.put`, a `gin.putValue`, or a standalone `gin.signal` emitted two system-scope fences per peer: one in `fenceBeforeSignal` and one in `signalPeer`. `signalPeer` now emits a single system-scope release immediately before the remote atomic, rather than an agent-scope release before it and a system-scope fence after it, which also stops payload stores from passing the signal cell. A clean-queue IPC `SignalInc` skips `fenceBeforeSignal` entirely, since it would only repeat that release. The SDMA `quiet()` before a signal now runs only when the call itself issued SDMA; previously any call carrying a signal or counter quieted the peer queue, including an IPC put after an earlier SDMA put to the same peer, which now relies on `Flush` to drain that queue. The `NCCL_GIN_ANVIL_SDMA_IPC_AGENT_FENCE` debug knob is removed. Each `Flush` thread drains the peers it owns and then clears exactly the bits it drained, over a mask disjoint from every other thread's. Previously `Flush` loaded a snapshot, quieted it, and cleared the whole word on the coop's leading thread, on every rank, which could drop a concurrent mark whose doorbell had not yet fired and then skip the quiet that mark needed. Clearing after the drain also keeps the invariant that a clear bit means a drained queue, so a second coop flushing the same context cannot return from `Flush` while the first is still inside `quiet()`. - #5708 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.30.7 for ROCm 10.1.0` section: > * `hipcc` is no longer accepted as the CMake CXX compiler for RCCL builds. Use `amdclang++` directly (`-DCMAKE_CXX_COMPILER=amdclang++` or set `ROCM_PATH`). `hipcc` was a wrapper around `amdclang++` that is being phased out across the ROCm stack. Only the CHANGELOG is changed. Version macros are left to the component owners. JIRA ID : AIROCDOC-4467 --------- Co-authored-by: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.com>

isaki001 requested review from a team and a lite review from Copilot September 28, 2026 17:25

github-actions Bot added the project: rccl label Sep 28, 2026

Copilot started reviewing on behalf of isaki001 September 28, 2026 17:26

View session

Copilot AI reviewed Sep 28, 2026

View reviewed changes

systems-assistant Bot added the organization: ROCm label Sep 28, 2026

amd-argus Bot reviewed Sep 28, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

Comment thread projects/rccl/src/include/ce_coll.h Resolved

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

github-actions Bot added the documentation Improvements or additions to documentation label Sep 29, 2026

amd-argus Bot reviewed Sep 29, 2026

View reviewed changes

isaki001 force-pushed the users/isaki001/tuning_infra_improvement branch from 9c4b828 to 91b4186 Compare October 5, 2026 21:42

amd-argus Bot reviewed Oct 5, 2026

View reviewed changes

amd-argus Bot reviewed Oct 6, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_arch_thresholds.cc Resolved

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

Comment thread projects/rccl/CHANGELOG.md Resolved

Comment thread projects/rccl/test/host/wrap-test.cc Resolved

amd-argus Bot reviewed Oct 6, 2026

View reviewed changes

amd-argus Bot reviewed Oct 7, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

amd-argus Bot reviewed Oct 7, 2026

View reviewed changes

Comment thread projects/rccl/test/host/wrap-test.cc Resolved

amd-argus Bot reviewed Oct 7, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_arch_thresholds.cc Resolved

amd-argus Bot reviewed Oct 9, 2026

View reviewed changes

isaki001 force-pushed the users/isaki001/tuning_infra_improvement branch from 61668c3 to a33f96c Compare October 9, 2026 17:05

amd-argus Bot reviewed Oct 9, 2026

View reviewed changes

Comment thread projects/rccl/test/CeAllReduceEligibilityTests.cpp Resolved

mustafabar reviewed Oct 9, 2026

View reviewed changes

pvallem reviewed Oct 12, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

pvallem reviewed Oct 12, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

pvallem reviewed Oct 12, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

pvallem reviewed Oct 12, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_arch_thresholds.cc Resolved

isaki001 added 10 commits October 19, 2026 14:20

fix test by setting cta policy and threshold

0e98c05

restore DDA LL128 A2A threshold to 1MiB

0896850

add back .forceCeScratch that was removed from formatting change

695cb4b

add missing type check for CE AR in ncclCeAvailable()

b442cb7

remove improper check non-registered CE AR, adjust test enforcement o…

03364e8

…f cta policy and force flag

update stale comment and fix improper test condition

1dd8bbc

adjust test to reflect DDA LL128 non-zero threshold

479f920

fix type check

2a906a8

address tests to reflect usage of RCCL_FORCE_CE

c9e65be

test adjustment for to accomodate increase of RCCL_CE_AR_STAGING_BYTES

8f97dd6

isaki001 force-pushed the users/isaki001/tuning_infra_improvement branch from 8f8b32d to 8f97dd6 Compare October 19, 2026 14:24

mustafabar approved these changes Oct 19, 2026

View reviewed changes

amd-argus Bot reviewed Oct 19, 2026

View reviewed changes

pvallem approved these changes Oct 19, 2026

View reviewed changes

pvallem reviewed Oct 19, 2026

View reviewed changes

Comment thread projects/rccl/src/enqueue/enqueue.cc Resolved

pvallem reviewed Oct 19, 2026

View reviewed changes

Comment thread projects/rccl/src/rccl_wrap.cc Resolved

sync CE selector for backup path for AG with rcclSelectAllGather

3dd82a0

amd-argus Bot reviewed Oct 19, 2026

View reviewed changes

isaki001 merged commit d0e8834 into develop Oct 20, 2026
74 of 81 checks passed

isaki001 deleted the users/isaki001/tuning_infra_improvement branch October 20, 2026 13:54

aduforest mentioned this pull request Oct 20, 2026

test(rccl): follow #12386's LL128 gating in the DDA launcher microtests #13819

Merged

atulkulk mentioned this pull request Oct 20, 2026

fix(rccl): honor abortFlag in DDA waits so revoke can complete #11245

Merged

1 task done

This was referenced Oct 21, 2026

[rocjitsu] Correct legacy floating min/max and median selection #13898

Merged

[rocjitsu] Expose backed CDNA3 and CDNA4 hardware identity fields #13899

Merged

[rocjitsu] Implement scratch LDS transfers and scalar cache discard #13900

Merged

pmoutsias-amd mentioned this pull request Nov 3, 2026

Remove duplicate RCCL 2.31.2 section and move post-branch entries to the RCCL 2.32.3 for ROCm 10.3.0 (Unreleased) section #14929

Merged

Conversation

isaki001 commented Sep 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

DDA LL128: Removed rcclParamDdaLL() == 0 guards inside ddaLL128ArOneShotEligible

AllGather dispatch non-registered CE branch fix

AllToAll CE registered guard: Added a2aCeRegWindow upper-bound check for the AllToAll CE registered path, preventing CE from being selected at message sizes beyond the arch-table ceiling.

ncclCeAvailable now has Float8 and ncclFloat8e5m2 checks for AllReduce

Test Changes

Issue Tracking

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Sep 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Policy Checks Passed

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot commented Sep 28, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

github-actions Bot commented Oct 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

isaki001 commented Sep 28, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 28, 2026 •

edited

Loading

github-actions Bot commented Oct 5, 2026 •

edited

Loading