feat(rccl): enable gfx1250 TDM symmetric kernels by default by alex-breslow-amd · Pull Request #13357 · ROCm/rocm-systems

alex-breslow-amd · 2026-10-13T15:08:48Z

Motivation

NCCL 2.31.2 release note: "On Blackwell, TMA kernels are now enabled by default and integrated in the cost model for symmetric registered memory."

RCCL's equivalents of those kernels use the gfx1250 Tensor Data Mover (TDM) and were added in #11850. They are AllGather_TmaST, AllReduce_RSxTmaLD_AGxTmaST and ReduceScatter_TmaLD. The symmetric cost model already includes their latency and bandwidth terms. However, NCCL_SYM_TMA_ENABLE still defaulted to 0, so users had to opt in. This PR makes them on by default, as NCCL does, and adds tests and documentation.

Technical Details

src/sym_kernels.cc: NCCL_SYM_TMA_ENABLE now defaults to 1. ncclSymkTmaAvailable() still allows these kernels only on gfx1250 builds whose SDK ships hip/amd_detail/amd_gfx1250_TDM.h. Other architectures don't change. Setting 0 restores the previous behavior, and 2 still forces the TDM kernels for A/B measurement.
test/host/sym-model-test.cc (new): host microtests that run the real symmetric cost model (src/tuning/sym_model.cc and sym_model/*.cc), which had no tests before. They are linked into rccl-UnitTestsMicroSymKernels. They check that:
- the default makes the TDM kernels available without forcing them;
- with the tuner's minimum selection time, the TDM kernels win only for large messages, and launch at 64 blocks of 16 warps;
- near the deep-loop size bar, the block count is reduced rather than the kernel being dropped;
- small sizes, misaligned buffers, groups of small works, unregistered windows, non-gfx1250 architectures and NCCL_SYM_TMA_ENABLE=0 all keep the TDM kernels out.
test/host/sym-kernels-test.cc: adds kernel-mask cases at the default parameter value, with no parameter hook installed.
-A 1 reporting (src/rccl_wrap.cc, src/include/rccl_common.h): rcclGetCollImplInfo and rcclSymKGetInfo now report a symmetric dispatch that resolves to a TDM kernel as SymTDM (new RCCL_SYMMETRIC_TDM), so rccl-tests -A 1 prints SymTDM instead of SYM for those sizes. The new value is for reporting only. Live dispatch still tags RCCL_SYMMETRIC, which the enqueue extraction gate and the dispatch switches compare against. The reporting query also never set symAligned16B, so the kernel mask always dropped the TDM kernels from the report. It now computes the bit from the window offsets, as ncclPrepareTasks does.
Rank-uniform TDM eligibility (src/include/enqueue.h, task_pretuning.cc, symmetric_sched.cc, rccl_wrap.cc): symAligned16B compared each rank's own window offsets. An in-place AllGather or ReduceScatter shifts those offsets by rank * perRankBytes, so with a per-rank slice that isn't a multiple of 16 B, ranks could pick different symmetric kernels and grids. The new ncclSymAligned16B also requires the per-rank slice to be 16 B aligned in that case. All three paths that fill the tuning input use it.
Docs: adds a "Tensor Data Mover symmetric kernels on gfx1250" section to rccl-usage-tips.rst and an NCCL_SYM_TMA_ENABLE entry to env-variables.rst. Also adds a CHANGELOG entry and updates the 2.31 sync triage row from "Skip" to "Adapt".

Known limitation (not changed here): the gfx1250 entries in rccl_arch_thresholds.cc run before the cost model. They keep registered AllReduce off the symmetric path entirely (symMaxR2[AR] = 0). They also allow symmetric AllGather only from 128 KiB to 2 MiB total, where the TDM kernel never wins. So with the default table, only ReduceScatter uses TDM. These thresholds were tuned before TDM existed and need retuning on an idle node; this should be tracked as a separate follow-up.

Issue Tracking

JIRA ID: AICOMRCCL-2008 (NCCL 2.31.2 release-note line item). A follow-up is needed to retune the gfx1250 symMaxR2 thresholds for the TDM kernels.

Test Plan

Host microtests: build and run rccl-UnitTestsMicro, rccl-UnitTestsMicroSymKernels, rccl-UnitTestsMicroTaskPrep, rccl-UnitTestsMicroEnqueue, rccl-UnitTestsMicroInit and rccl-UnitTestsMicroCe.
Host microtests for the in-place alignment rule: SchedulerMicrotest.SymBatchAligned16B_InPlace* in rccl-UnitTestsMicroWarpSpeed.
Hardware: run rccl-tests on 4× gfx1250 with symmetric registration (-R 2 -c 1, single process with -g 4), from 64 KiB to 1 GiB. Compare NCCL_SYM_TMA_ENABLE unset, 0, 1 and 2, with and without RCCL_IGNORE_ARCH_TABLE=1. Use NCCL_DEBUG_SUBSYS=TUNING to record which kernel is selected.
Forced-TDM correctness: run with NCCL_SYM_TMA_ENABLE=2 RCCL_IGNORE_ARCH_TABLE=1 for AllReduce and ReduceScatter sum with float, half, bfloat16, fp8_e4m3 and fp8_e5m2, and AllGather with float, bfloat16 and int8. Repeat at non-power-of-two sizes that are still 16 B-aligned per rank.

Test Result

Host microtests: all 3,496 pass across rccl-UnitTestsMicro, -Ce, -Enqueue, -Init, -SymKernels (67), -TaskPrep, -WarpSpeed (the scheduler tests), -Diagnostics and -GinHost. None were skipped, so the TDM toolchain was detected.
-A 1 label: reduce_scatter_perf -g 4 -b 64M -e 1G -f 4 -R 2 -A 1 prints SYM at 64 MiB and SymTDM at 256 MiB and 1 GiB. With RCCL_IGNORE_ARCH_TABLE=1, every SYM/SymTDM label for AllReduce, AllGather and ReduceScatter matches the kernel in the [Symmetric] log line. The CollImplInfo GPU tests that compare the reported algorithm with the dispatch log pass for all three symmetric collectives. CollImplInfo.AllReduceCeRegisteredMatchesDispatchLog timed out in the debug build. That test runs the Copy Engine path, which never reaches the symmetric query (the test notes a known hang, ROCM-32044).
Correctness: 0 wrong elements in every hardware run, both forced and tuner-selected.
Kernel selection: with the variable unset, the tuner picks TDM for AllReduce from 16 MiB, AllGather from 64 MiB and ReduceScatter from 256 MiB. Below those sizes it keeps the non-TDM kernels.
Bandwidth: bus bandwidth at 1 GiB with the arch table bypassed, normalized to the non-TDM AllGather kernel (AllGather_ST) at the same size, which is 1.00×. The machine was shared with another GPU workload, so only the ratios are meaningful.

Collective Non-TDM TDM

AllReduce 1.60× 8.03–10.56×

AllGather 1.00× (baseline) 6.44–6.72×

ReduceScatter 2.58–2.64× 5.23–5.72×

With the default arch table, ReduceScatter now selects ReduceScatter_TmaLD at 256 MiB and above. At 1 GiB it reaches 5.23–6.75× the baseline, up from 2.56–2.64× before.
Not covered:
- rccl-tests prints no result rows for -o avg on this build, with or without TDM, so ReduceScatter avg with TDM is covered only by the host tests.
- The MPI-based symmetric-kernel tests were not run, because MPI isn't available on the test machine.
- The rank-uniform alignment fix is covered by host tests only. A hardware run of an in-place, misaligned-slice ReduceScatter stalled before its first collective was logged, and the node then stopped completing even small collectives, so there is no hardware result for it yet.

Submission Checklist

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

therock-pr-bot · 2026-10-13T15:09:41Z

✅ 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/enqueue/task_prep/task_pretuning.cc`, `projects/rccl/src/include/enqueue.h`, `projects/rccl/src/include/rccl_common.h`, `projects/rccl/src/include/sym_kernels.h`, `projects/rccl/src/rccl_wrap.cc` (+19 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?

alex-breslow-amd · 2026-10-13T15:26:02Z

@copilot resolve the merge conflicts in this pull request

Copilot

🟡 Changes recommended

The documentation misstates supported datatypes and window-registration requirements, with corresponding registration coverage missing from the new tests.

4 open findings

What changed in this PR

Enables gfx1250 TDM symmetric kernels by default and documents/tests their cost-model integration.

Changes:

Defaults NCCL_SYM_TMA_ENABLE to 1.
Adds kernel-mask and cost-model tests.
Documents availability, tuning, and architecture constraints.

File	Description
`projects/rccl/src/sym_kernels.cc`	Enables TDM kernels by default.
`projects/rccl/test/host/sym-model-test.cc`	Adds symmetric cost-model tests.
`projects/rccl/test/host/sym-kernels-test.cc`	Tests default kernel masks.
`projects/rccl/test/host/CMakeLists.txt`	Builds the new model tests.
`projects/rccl/docs/how-to/rccl-usage-tips.rst`	Adds TDM usage guidance.
`projects/rccl/docs/api-reference/env-variables.rst`	Documents the environment variable.
`projects/rccl/docs/dev/nccl-2.31-sync-triage.md`	Updates NCCL adaptation status.
`projects/rccl/CHANGELOG.md`	Records default enablement.

🧠 Review effort: Balanced

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

Copilot · 2026-10-13T15:36:14Z

@copilot resolve the merge conflicts in this pull request

Merged develop and resolved the changelog conflict, preserving both entries. Merge commit: 83eb6030a.

Kapil-Shyam-Pawar

As discussed, I feel like it would be better if we include the SYM-TDM getting printed in the algo column with this change
without that it would be very difficult to identify which msg ran with TDM without having the tuning logs

amd-argus

ARGUS

Findings

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

projects/rccl/src/sym_kernels.cc:390 (High) symAligned16B is per-rank and now decides kernel choice, so ranks can launch different kernels and grids and hang symAligned16B becomes a live kernel-selection input at sym_kernels.cc:390 once the default is 1, and it is a per-rank value: symBatchAligned16B() (symmetric_sched.cc:75-83) tests uint32_t(inputOff - outputOff) % 16, which for an in-place AllGather or ReduceScatter is rank * chunkBytes. ''' in-place ReduceScatter, fp32, 4 ranks, recvcount 67108865, chunkBytes 268435460 rank 0 0 * chunkBytes % 16 = 0 aligned -> ReduceScatter_TmaLD, 64 blocks rank 1 1 * chunkBytes % 16 = 4 not -> ReduceScatter_LD, CTA-search count rank 2 2 * chunkBytes % 16 = 8 not -> ReduceScatter_LD rank 3 3 * chunkBytes % 16 = 12 not -> ReduceScatter_LD ''' Nothing reconciles devFuncId or nMaxChannels across ranks (symmetric_sched.cc:271-273), and collectives.cc:319 says cross-rank agreement belongs in ncclMakeSymmetricTaskList, where it was never added. Since the LSA barrier is indexed by blockIdx.x (reduce_scatter.cuh:359) over lsaBarrierCount = ncclSymkMaxBlocks, rank 0's surplus blocks wait on peers that never launch them. NCCL_CHECK_MODE=DEBUG_GLOBAL does not catch it: argcheck.cc:130 and :154 exempt exactly the side that varies per rank for these two collectives. In-place AllReduce is immune, since the offsets are equal on every rank. Could the TDM gate key only on communicator-uniform facts?
projects/rccl/CHANGELOG.md:77 (Low) The 2.31.2 entry still calls the TDM kernels opt-in, contradicting the new 2.32.3 entry in the same shipped changelog CHANGELOG.md:77, in the RCCL 2.31.2 for ROCm 10.2.0 (Unreleased) section, still says the TDM kernels are Opt-in via NCCL_SYM_TMA_ENABLE=1, which the new bullet at line 36 now contradicts. Both section headers say Unreleased for the same ROCm 10.2.0, so both ship together and a reader of the 2.31.2 entry alone is told the feature is off by default. Could line 77 be reworded to point at the 2.32.3 change?

Bottom line: Worth settling before approval: symAligned16B is a per-rank value and now decides kernel choice, so an in-place AllGather or ReduceScatter whose per-rank chunk is not 16-byte aligned can split the communicator and hang.

_{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/src/sym_kernels.cc:390 (High) symAligned16B is per-rank and now decides kernel choice, so ranks can launch different kernels and grids and hang symAligned16B is a per-rank value that now decides kernel choice at sym_kernels.cc:390, which the default of 1 makes live for the first time. symBatchAligned16B (symmetric_sched.cc:75-83) tests uint32_t(inputOff - outputOff) % 16, and for an in-place AllGather or ReduceScatter that difference is rank * chunkBytes. ''' in-place ReduceScatter, fp32, 4 ranks, recvcount 67108865, chunkBytes 268435460 rank 0 0 * chunkBytes % 16 = 0 aligned -> ReduceScatter_TmaLD, 64 blocks rank 1 1 * chunkBytes % 16 = 4 not -> ReduceScatter_LD, CTA-search count rank 2 2 * chunkBytes % 16 = 8 not -> ReduceScatter_LD rank 3 3 * chunkBytes % 16 = 12 not -> ReduceScatter_LD ''' Nothing reconciles devFuncId or nMaxChannels across ranks (symmetric_sched.cc:271-273), and collectives.cc:319 says cross-rank agreement belongs in ncclMakeSymmetricTaskList, where it was never added. The LSA barrier is indexed by blockIdx.x (reduce_scatter.cuh:359) over lsaBarrierCount = ncclSymkMaxBlocks, so rank 0's surplus blocks wait on peers that never launch them. NCCL_CHECK_MODE=DEBUG_GLOBAL does not catch it: argcheck.cc:130 and :154 exempt exactly the side that varies per rank for these two collectives, and it is off by default. In-place AllReduce is immune, since the offsets are equal on every rank. Could the TDM gate key only on communicator-uniform facts?

Bottom line: Worth settling before approval: symAligned16B is still per-rank and now decides kernel choice, so an in-place AllGather or ReduceScatter whose per-rank chunk is not 16-byte aligned can split the communicator and hang.

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

alex-breslow-amd · 2026-10-13T18:36:49Z

@amd-argus, on the two findings in your review summary:

Per-rank symAligned16B (High): Agreed, and fixed in 266c35c. Upstream NCCL 2.32.3 computes the bit the same way, but with TDM on by default RCCL can now reach the split. The new ncclSymAligned16B helper (src/include/enqueue.h) keeps the existing window-offset check. For in-place AllGather and ReduceScatter it also requires the per-rank slice to be a multiple of 16 bytes, so every rank gets the same answer. The device kernels' own per-rank deep-tier gate is unchanged: a misaligned rank inside a TDM kernel still falls back to the vector path within the same kernel and grid. Pretuning, symBatchAligned16B and the rcclSymkQuery reporting path all use the helper.

New host tests in scheduler-test.cc check all four ranks for in-place AllGather and ReduceScatter:

with a 24-byte slice, every rank gets false (the local check alone passes on ranks 0 and 2);
with a 32-byte slice, every rank gets true;
out-of-place with an odd slice is unaffected.

wrap-test.cc covers the same case through the reporting query. I couldn't confirm the hang on hardware: the reproduction attempt stalled before the first collective was logged, and the node then needed attention.

CHANGELOG 2.31.2 entry (Low): Fixed in 266c35c. The entry now points at the 2.32.3 default change instead of saying "opt-in".

amd-argus

ARGUS

Bottom line: Worth settling before approval: the gfx1250 TDM kernels are still on by default with no opt-in, and the new rank-uniform alignment bit is only tested at rank 0, where rank * perRankBytes is zero.

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

amd-argus

ARGUS

This PR changes 1,078 lines across 21 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: Worth settling before approval: NCCL_SYM_TMA_ENABLE still defaults to 1 at sym_kernels.cc:146, so the gfx1250 TDM kernels ship tuner-selectable with no opt-in.

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

…#13787) The tests in #13357 miss AllGather/ReduceScatter SymTDM reporting, the queried collective, the byte-sized slice and the full LL/TDM kernel sets, and its in-place GPU tests run in no config. Fix: add those host tests, make the GPU tests check the kernel that ran, and list them in the MI455 configs. JIRA ID: AICOMRCCL-2008

amd-argus

ARGUS

This PR changes 1,463 lines across 25 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: Nothing new worth raising; still open from earlier: sym_kernels.cc:146 and 1 more. Worth clearing before approval.

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

NCCL 2.31.2 enables its TMA symmetric kernels by default on Blackwell and costs them in the symmetric cost model. RCCL's counterparts stage through the gfx1250 Tensor Data Mover and were already costed, but were opt-in. - NCCL_SYM_TMA_ENABLE now defaults to 1. ncclSymkTmaAvailable() still limits these kernels to gfx1250 builds whose SDK ships the TDM descriptor header, so other architectures are unchanged. - Add host microtests for the real symmetric cost model (sym-model-test.cc). They check that the default leaves the choice to the model, that Tma kernels win only for large messages and launch at the widest block count with 16 warps, and that the deep-loop, alignment, grouping, registration and arch gates keep them out elsewhere. - Add default-parameter mask tests to sym-kernels-test.cc. - Document NCCL_SYM_TMA_ENABLE and the TDM kernels in the env-variable reference and the usage tips, and add a CHANGELOG entry. Co-authored-by: Cursor <cursoragent@cursor.com>

…lInfo rccl-tests -A 1 labelled every symmetric dispatch SYM, so it could not show when the gfx1250 Tensor Data Mover kernels ran. Add a reporting-only RCCL_SYMMETRIC_TDM ("SymTDM") that rcclSymkQuery returns when the cost model picks a Tma kernel; live dispatch still tags RCCL_SYMMETRIC, which the enqueue extraction gate and dispatch switches compare against. The query also left symAligned16B unset, so the mask always dropped the TDM kernels and the report could disagree with the dispatched kernel. Compute it from the window offsets as ncclPrepareTasks does; null buffers (the rcclSymKGetInfo ABI) count as aligned. Co-authored-by: Cursor <cursoragent@cursor.com>

symAligned16B compared each rank's own input/output window offsets, which an in-place AllGather or ReduceScatter shifts by rank * perRankBytes. With a per-rank slice that is not a multiple of 16 B, ranks reached different answers, so with NCCL_SYM_TMA_ENABLE now defaulting to 1 some ranks could launch a Tma kernel at 64 blocks while others launched the vector kernel with fewer, and the surplus blocks would wait at the LSA barrier. ncclSymAligned16B also requires the per-rank slice to be 16 B aligned for in-place AllGather/ReduceScatter, and the pretuning, symmetric scheduler and rcclSymkQuery paths all use it. rcclSymkQuery now takes the alignment bit from its callers, which already hold both windows, instead of looking them up again. Tests: rank-uniform in-place cases for symBatchAligned16B and the reporting query; direct rcclSymkKernelIdIsTma/IsLL coverage; a default-value test that runs without the TDM header; one-sided registration cases for the model; skip guards where the gate under test would not run; wave64 for GFX9 archs in the model fixture; a shared vector-kernel helper; and a Sym-prefixed suite name. Drop two default-param mask tests that the old default also passed. Docs: drop NVIDIA hardware comparisons, describe the in-place alignment rule, allow any datatype for AllGather, remove the arch-table paragraph, reframe the 2.31 triage row around behavior, and fix the 2.31.2 changelog entry that still called the TDM kernels opt-in. Co-authored-by: Cursor <cursoragent@cursor.com>

…n in-place slices In-place AllGather and ReduceScatter on symmetric windows with 4 MiB slices, one 16 B aligned and one 12 B off. NCCL_SYM_TMA_ENABLE=2 reaches the TDM kernels on gfx1250; each case runs in an isolated child with a timeout, so a rank-divergent kernel choice fails instead of hanging. Named under the ReduceScatter/AllGather suites so the standard CI tiers pick them up. On 4x gfx1250 the pre-fix library times out on ReduceScatterUnalignedSlice (rank 0 ReduceScatter_TmaLD vs ReduceScatter_LD elsewhere); the fixed library passes all four cases. Co-authored-by: Cursor <cursoragent@cursor.com>

Move the in-place ReduceScatter selector test to rank 2 of 4 and add the AllGather counterpart. With a 24 B slice, rank 2's operands differ by 48 B, which passes the plain modulo-16 check, so only the rank-aware in-place rule reports them unaligned. Replacing comm->rank with 0 at the three rcclSelect* call sites now fails both tests. Co-authored-by: Cursor <cursoragent@cursor.com>

…#13787) The tests in #13357 miss AllGather/ReduceScatter SymTDM reporting, the queried collective, the byte-sized slice and the full LL/TDM kernel sets, and its in-place GPU tests run in no config. Fix: add those host tests, make the GPU tests check the kernel that ran, and list them in the MI455 configs. JIRA ID: AICOMRCCL-2008

amd-argus

ARGUS

This PR changes 1,463 lines across 25 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: Nothing new worth raising; still open from earlier: sym_kernels.cc:148 and 1 more. Worth clearing before approval.

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

alex-breslow-amd · 2026-10-21T21:46:49Z

@Kapil-Shyam-Pawar, the SymTDM label you asked for is in 031e03f: rccl-tests -A 1 now prints SymTDM when a symmetric dispatch picks a TDM kernel.

The branch was rebased on Oct 20, so the commit hashes in earlier replies changed. Same content, new hashes:

5da9b4c → 031e03f (SymTDM reporting; the alignment bit in the reporting query)
266c35c → 12ae942 (rank-uniform alignment fix, CHANGELOG, docs and most review fixes)
2d7cc37 → 28b61c8 (tests at a non-zero rank)
662f3fc → be48565 (test(rccl): close test gaps in the gfx1250 TDM symmetric kernel tests #13787 test gaps)

The GPU regression test for the in-place fix is 9f5752e.

ncclSymAligned16B fell back to the raw address of a buffer outside any window. Raw addresses differ by rank, so a send-only or recv-only registration (ncclReduceScatterConfig reaches the symmetric path without going through rcclSelectReduceScatter) could make ranks disagree on symAligned16B, pick different symmetric kernels, and hang at the LSA barrier. Both windows must now exist and carry NCCL_WIN_COLL_SYMMETRIC. Tests: scheduler, task pretuning and wrap tests cover missing and non-symmetric windows, in-place slices through fillCollTuningInput, and record symAligned16B without static state. Docs say NCCL_SYM_TMA_ENABLE must match on every rank. Co-authored-by: Cursor <cursoragent@cursor.com>

- SymTdmTestHelpers.h holds SymTmaEnable, the arch setters and TmaReachableOnThisBuild for sym-kernels-test and sym-model-test. - sym-model-test: SymModelTdmReachableTest skips in SetUp instead of per-test guards; uses rcclSymkKernelIdIsTma directly. - ncclTeamRail lives once in devcomm_fakes.cc. - ncclSymBufAllocPerRank allocates and registers per-rank symmetric buffers for CollImplInfoTests and SymmetricInPlaceSliceTests, and reports a registration that returns no window as a failure. - IsSymmetricAlgoName replaces the duplicated "SYM"/"SymTDM" checks. Co-authored-by: Cursor <cursoragent@cursor.com>

The receive buffer is now rank 2's slice of the send buffer, so the task is in place at rank 2 and out of place at rank 1. An in-place check that ignored the rank would give the same answer at both. Co-authored-by: Cursor <cursoragent@cursor.com>

…metricMemoryTDMEnable Conflict in CHANGELOG.md: kept both sides' bullets. wrap-test.cc: a test from develop calls SelectSymkTuning without the func argument this branch added; it now passes ncclFuncAllReduce. Co-authored-by: Cursor <cursoragent@cursor.com>

amd-argus

ARGUS

This PR changes 1,843 lines across 32 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: Nothing new worth raising; still open from earlier: sym_kernels.cc:148. Worth clearing before approval.

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

alex-breslow-amd · 2026-10-22T02:17:04Z

ARGUS

This PR changes 1,843 lines across 32 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: Nothing new worth raising; still open from earlier: sym_kernels.cc:148. Worth clearing before approval.

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

This PR is only as long as it is to satisfy all of the Argus nitpicks and testing.

systems-assistant · 2026-10-26T06:03:39Z

TheRock Submodule Bump Activity

Newest first

2026-10-26 — Included in TheRock via ROCm/TheRock#9480.

…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>

alex-breslow-amd added the project: rccl label Oct 13, 2026

github-actions Bot added the documentation Improvements or additions to documentation label Oct 13, 2026

systems-assistant Bot added the organization: ROCm label Oct 13, 2026

alex-breslow-amd marked this pull request as ready for review October 13, 2026 15:25

alex-breslow-amd requested review from a team and a balanced review from Copilot October 13, 2026 15:25

Copilot started reviewing on behalf of alex-breslow-amd October 13, 2026 15:27

View session

Copilot AI reviewed Oct 13, 2026

View reviewed changes

Comment thread projects/rccl/test/host/sym-model-test.cc Resolved

Comment thread projects/rccl/docs/api-reference/env-variables.rst Resolved

Comment thread projects/rccl/docs/how-to/rccl-usage-tips.rst Resolved

Comment thread projects/rccl/docs/how-to/rccl-usage-tips.rst Resolved

Copilot started work on behalf of alex-breslow-amd October 13, 2026 15:32

View session

Copilot finished work on behalf of alex-breslow-amd October 13, 2026 15:38

Kapil-Shyam-Pawar reviewed Oct 13, 2026

View reviewed changes

amd-argus Bot reviewed Oct 13, 2026

View reviewed changes

garzaran reviewed Oct 13, 2026

View reviewed changes

Comment thread projects/rccl/docs/api-reference/env-variables.rst Resolved

Comment thread projects/rccl/docs/dev/nccl-2.31-sync-triage.md Resolved

Comment thread projects/rccl/docs/how-to/rccl-usage-tips.rst Resolved

Comment thread projects/rccl/docs/how-to/rccl-usage-tips.rst Resolved

amd-argus Bot reviewed Oct 13, 2026

View reviewed changes

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

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

amd-argus Bot reviewed Oct 13, 2026

View reviewed changes

amd-argus Bot reviewed Oct 14, 2026

View reviewed changes

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

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

pvallem reviewed Oct 19, 2026

View reviewed changes

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

pvallem reviewed Oct 19, 2026

View reviewed changes

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

pvallem reviewed Oct 19, 2026

View reviewed changes

Comment thread projects/rccl/test/host/sym-model-test.cc Resolved

pvallem reviewed Oct 19, 2026

View reviewed changes

Comment thread projects/rccl/test/host/sym-model-test.cc Resolved

pvallem reviewed Oct 19, 2026

View reviewed changes

Comment thread projects/rccl/test/host/sym-model-test.cc Resolved

pvallem reviewed Oct 19, 2026

View reviewed changes

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

pvallem reviewed Oct 19, 2026

View reviewed changes

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

pvallem mentioned this pull request Oct 20, 2026

test(rccl): close test gaps in the gfx1250 TDM symmetric kernel tests #13787

Merged

amd-argus Bot reviewed Oct 20, 2026

View reviewed changes

Alex Breslow and others added 6 commits October 20, 2026 22:59

alex-breslow-amd force-pushed the users/abreslow/symmetricMemoryTDMEnable branch from 662f3fc to be48565 Compare October 20, 2026 23:03

amd-argus Bot reviewed Oct 20, 2026

View reviewed changes

Alex Breslow and others added 4 commits October 21, 2026 22:18

Kapil-Shyam-Pawar approved these changes Oct 21, 2026

View reviewed changes

amd-argus Bot reviewed Oct 21, 2026

View reviewed changes

pvallem approved these changes Oct 22, 2026

View reviewed changes

alex-breslow-amd merged commit a08bc31 into develop Oct 23, 2026
109 of 115 checks passed

alex-breslow-amd deleted the users/abreslow/symmetricMemoryTDMEnable branch October 23, 2026 06:42

pvallem mentioned this pull request Oct 23, 2026

fix(rccl): Enforce Zero CU Policy and Improve Handling of Forcing Extended Algos #13098

Merged

1 task

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

Collective	Non-TDM	TDM
AllReduce	1.60×	8.03–10.56×
AllGather	1.00× (baseline)	6.44–6.72×
ReduceScatter	2.58–2.64×	5.23–5.72×

Conversation

alex-breslow-amd commented Oct 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Issue Tracking

Test Plan

Test Result

Submission Checklist

Uh oh!

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

Uh oh!

✅ All Policy Checks Passed

Uh oh!

alex-breslow-amd commented Oct 13, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI commented Oct 13, 2026

Uh oh!

Kapil-Shyam-Pawar left a comment

Choose a reason for hiding this comment

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!

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!

alex-breslow-amd commented Oct 13, 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!

Uh oh!

Uh oh!

Uh oh!

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

alex-breslow-amd commented Oct 13, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 13, 2026 •

edited

Loading