feat(rccl): stage symmetric kernel tiles with gfx1250 TDM instead of TMA by alex-breslow-amd · Pull Request #11850 · ROCm/rocm-systems

alex-breslow-amd · 2026-09-18T14:04:03Z

Overview

Ports TMA calls in the symmetric memory kernels to use AMD's tensor data mover. Fixes some of the tuning logic so that up to 1/4 of the WGPs on a GPU can be used. Without this, only at most 14 WGPs were observed to be used. Raises performance significantly for large data sizes.

The code is gated by the NCCL_SYM_TMA_ENABLE=1 environment variable and only applies to when symmetric memory is opted in by the application.

Some Details

gfx1250 is wave32, so NVIDIA's tile geometry carries over unchanged: the 256 B-aligned AllGather tier is 16 packs x 32 lanes x 16 B = 8 KiB, the reduce tiers are 2 x 8 x 32 x 16 B = 8 KiB, both inside the existing 10 KiB/warp scratch. Static LDS is 0 (these kernels are entirely dynamic-LDS); at 256 threads that is 8 warps x 10 KiB = 80 KiB of MI450's 320 KiB. No resizing needed. Further work will look at tuning this given that more LDS would allow for more in-flight work to hide the latency.

Issue Tracking

JIRA ID: AICOMRCCL-1700 and AICOMRCCL-1203

Test Plan

Build a full gfx1250 librccl.so (ROCm 7.15) and confirm it links.
Confirm the ported path is genuinely live rather than silently falling back to the vector path, by unbundling the shipped gfx1250 code object and disassembling it for TDM opcodes.
Regression: compile the same TUs for a non-TDM arch (gfx950) and confirm zero TDM opcodes.
Regression: build on ROCm 7.0.2.1, which ships no TDM descriptor header at all, so TDM_SUPPORTED is 0 via __has_include.
Run the symmetric kernel-count leak test.

Test Result

1. Full gfx1250 build: links clean, 320/320 targets. Only pre-existing -Wpass-failed=transform-warning unroll warnings.
2. TDM opcodes in the shipped code object (unbundled hipv4-amdgcn-amd-amdhsa--gfx1250 from sym_*.o, disassembled):

Kernel	`tensor_load_to_lds`	`tensor_store_from_lds`	`s_wait_tensorcnt`
`AllGather_TmaST`	12	18	8
`ReduceScatter_TmaLD`	18	3	4
`AllReduce_RSxTmaLD_AGxTmaST`	18	9	4
Store counts track each collective's fan-out: AllGather writes every peer, ReduceScatter writes one local output, AllReduce writes every peer. In the disassembly the nine back-to-back `D0714001` (`tensor_store_from_lds`) followed by a single `BFCB0000` (`s_wait_tensorcnt 0`) is exactly the intended peer fan-out plus one drain.
3. gfx950: all three TUs compile, zero TDM opcodes — the arch gate holds even for the `EnableTma=true` kernels, which degrade to the vector path.
4. ROCm 7.0.2.1 / gfx950: host and device symmetric objects build; generator emits 42 kernels (no Tma), confirming no new toolchain requirement for existing arches.
5. Kernel-count leak test: 15/15 pass, including a new case pinning the gfx1250 baseline at 58 and asserting the arch algos are purely additive.

Not covered

Two known performance follow-ups, deliberately left out of scope:

tdmCopy.h's window is single-buffered with RAW/WAR waits, so tdmWait() after each stage serializes more than TMA's mbarrier pipeline did. Real overlap needs a second LDS window — there is ample room in 320 KiB.
Aligned=false leaves head/tail peel branches in place. At runtime a 256 B-aligned tile still issues one descriptor, but the 16 B-aligned tier issues up to three; plumbing Aligned through bcastDeep would tighten the 256 B tier.

Submission Checklist

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

therock-pr-bot · 2026-09-18T14:04:50Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 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?

therock-pr-bot · 2026-09-18T14:04:52Z

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

❌ PR Description

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

Copilot

Copilot review overview

🟡 Changes recommended

Unresolved critical TDM alignment and moderate gating/tuning issues block approval.

Get a fresh assessment by requesting another Copilot review.

Review effort: Lite
Findings: 1 · 2

Open (3)

TDM transfers can use mismatched global and LDS alignment · New
Gate TMA kernels on TDM toolchain availability · New
Even CTA rounding can violate the requested minimum · New

What changed in this PR

Ports RCCL symmetric kernels to gfx1250 TDM staging and expands launch tuning for better WGP utilization.

Changes:

Adds TDM tile staging to symmetric collectives.
Generates gfx1250-specific kernels with architecture gating.
Updates launch tuning and kernel-count regression coverage.

File	Changes and review status
`projects/rccl/test/kernel-count/tests/test_symmetric_kernels.py`	Adds gfx1250 kernel-count coverage.
`projects/rccl/src/tuning/sym_model.cc`	Adjusts TDM launch sizing. Moderate (3 votes): minimum CTA counts can be rounded below the requested floor.
`projects/rccl/src/sym_kernels.cc`	Adds gfx1250 runtime gating. Moderate (2 votes): also gate on toolchain TDM availability.
`projects/rccl/src/device/symmetric/reduce_scatter.cuh`	Integrates asynchronous tile staging.
`projects/rccl/src/device/symmetric/primitives.cuh`	Adds TDM staging wrappers. Critical (1 vote): global/LDS alignment may mismatch for the default path. Nit (1 vote): add gfx1250 runtime correctness coverage.
`projects/rccl/src/device/symmetric/generate.py`	Generates architecture-specific kernels.
`projects/rccl/src/device/symmetric/all_reduce.cuh`	Integrates asynchronous tile staging.
`projects/rccl/src/device/symmetric/all_gather.cuh`	Integrates asynchronous tile staging.
`projects/rccl/src/CMakeLists.txt`	Passes GPU targets to kernel generation.

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

amd-argus

ARGUS

Bottom line: Worth settling before approval: tdmBlocks() sizes the grid from per-rank comm->cuCount, so ranks with different CU counts can launch different grid widths and hang on the block-indexed barrier.

_{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/CHANGELOG.md:30 (Low) No changelog entry, and the 2.31.2 Known issues list still denies the feature this PR delivers CHANGELOG.md:30, under the 2.31.2 Known issues heading, still lists "TMA-based symmetric kernels" among the NVIDIA-only features unavailable in RCCL, which this PR delivers on gfx1250 behind NCCL_SYM_TMA_ENABLE. The comparable TDM SIMPLE path took an Added entry at CHANGELOG.md:40 despite also being opt-in and out of the default build. Drop the clause and add an Added entry under 2.31.2?

Bottom line: Worth settling before approval: BytePerChunk at all_reduce.cuh:345 sits outside the #if NCCL_SYMK_ASYNC_TILE that guards its UnrollPacks, so a non-async-tile pass reduces a quarter of the range it skips past. Same line at reduce_scatter.cuh:301.

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

mustafabar

I went over the PR with @alex-breslow-amd. I would vote it in as it helps us reduce the performance gap between symmetric and non-symmetric kernel on MI450 (e.g. traditional ring/tree and CE based collectives).
Please include in the Jira the A/B performance delta on MI450 MCH SPX and/or MI455 1p4g. Also, with latest develop compare with the OOB non-symmetric case (-R 0) and set RCCL_IGNORE_ARCH_TABLE=0

amd-argus

ARGUS

Bottom line: Worth settling before approval: nothing bounds ncclTmaShmemScratchWarpSize()'s 18 KiB per warp times 16 warps (288 KiB) against the dynamic LDS the launch is actually granted.

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

nusislam · 2026-09-23T22:03:19Z

There is no check that nWarps * 18KiB <= args->maxDynamicSmem. A static_assert on sizeof(tmaSmemStruct) vs 18 KiB, plus a device-side compare to maxDynamicSmem, would catch the next tile widening. The PR description still says “8 warps x 10KiB", no resize"; that is no longer true.

nusislam · 2026-09-23T22:04:43Z

This is a user-visible opt-in path, same class as SIMPLE TDM. Add a changelog entry and drop any remaining “TMA-based symmetric kernels are NVIDIA-only” known-issue wording on the target branch.

The deep loops in the symmetric AllGather/AllReduce/ReduceScatter kernels move a tile through per-warp shared memory with a DMA engine rather than per-lane vector loads. That path was NVIDIA TMA behind `#if __CUDA_ARCH__ >= 1000`, so it compiled out entirely on ROCm. Route it through gfx1250's Tensor Data Mover. `cuda::barrier<cuda::thread_scope_block>` has no HIP equivalent (ROCm libhipcxx lists it unsupported), and it is not what is needed anyway: the barrier was doing TMA transaction accounting, whereas TDM completion is a per-wave TENSORcnt. The two engines also differ in who issues a transfer. Rather than duplicate each async block per vendor, primitives.cuh gains a set of ncclSymkTile* wrappers that are called by the whole warp with wave-uniform arguments and narrow to a single thread internally where the hardware wants that, so the loop bodies keep one shape and all CUDA TMA API stays confined to that one file. Four differences needed more than a macro swap: * Issue is wave-collective. A TDM descriptor lives in SGPRs, so every lane calls in with identical arguments. * bcastDeep parked lanes != 0 out of the loop via `skip`, which would have left TDM with a partial wave. * bcastMultimem's cursor carries a per-lane slot. TMA read lane 0's copy as the tile base; that slot is now backed out explicitly to get a uniform base. * fence_proxy_async becomes __threadfence_block() to publish the lanes' reduced tile to the engine about to read it. Gating is on TDM_SUPPORTED (arch + builtin + SDK descriptor header), not ENABLE_TDM_SIMPLE, which is the SIMPLE-protocol path and reserves 160 KiB of static LDS these kernels do not want. Every other arch, and the host pass, keep the vector path untouched. Host side, ncclSymkTmaAvailable() gated on `minCompCap >= 100`, a CUDA compute capability AMD never reports, so it now keys on the arch carrying the mover. These kernels are all LSA-scoped, so every participating rank shares that arch. generate.py never emitted the Tma kernel symbols, so the device code had no instantiation and `if constexpr (EnableTma)` discarded it unconditionally. It now emits AllGather_TmaST, ReduceScatter_TmaLD and AllReduce_RSxTmaLD_AGxTmaST when gfx1250 is in GPU_TARGETS (42 -> 58 kernels), and nothing extra otherwise. TmaSTMC stays out: it needs multimem, which ROCm has no equivalent for. Still opt-in via NCCL_SYM_TMA_ENABLE=1. Co-authored-by: Cursor <cursoragent@cursor.com>

Reverts 20a9b0e. That commit split the tile cache policy into kNcclSymkTileLocalPolicy (NT/DEV) for this rank's own copy and kNcclSymkTilePeerPolicy (NT/SYS) for peers, threaded a CachePolicy template parameter through ncclSymkTileLoad/ncclSymkTileStore to carry it, and peeled the own-rank store out of the fan-out loop in bcastDeep and allreduceDeep so the remaining loop emits only one policy. Measured across 1 MB - 1 GB with the TDM kernel forced at every size, it does not pay for itself. AllReduce and ReduceScatter are unaffected: every size lands within 1-2%, and ReduceScatter is the strongest case since its only store is the local one, so the revert moves 100% of its store traffic from DEV back to SYS and still lands within 1% at all eleven sizes. AllGather does move, in both directions. Over seven interleaved repetitions it gains 3.5-6% at 32-64 MB without the split and loses 4-6% at 512 MB and 1 GB, both outside the run-to-run band. Removed regardless, because of where those effects land. The Copy Engine already owns the sizes device scope was helping: NCCL_SYM_CE_THRESHOLD is 8388608 elements of sendcount, which is 128 MB of total AllGather for f32 on 4 ranks, so 256 MB and up dispatch hipMemcpy and beat the TDM kernel by 20-33% anyway. The large-size win is a win on a path that does not run, while the 32-64 MB gain is exactly where the kernel has the most untapped headroom. Net: -73 lines, one cache policy instead of two, no template parameter on the tile wrappers, and no own-rank peel in two deep loops. Co-authored-by: Cursor <cursoragent@cursor.com>

amd-argus

ARGUS

Bottom line: Worth settling before approval: gfx1250 is in DEFAULT_GPUS, so have_tdm fires in a stock build and the 42-name stub list in test/host/fakes/sym_kernels_test_stubs.h leaves 32 undefined references, breaking the link of both microtest binaries.

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

sdarisala

The TDM changes in the PR is enabled for gfx1250 target only. Going forward, A0 silicon will start reporting as gfx1250-strict. The gfx1250 target detection checks should also work for gfx1250 (I think). The new test cases implemented should catch it, if it doesn't.
Will the new test cases be picked up by nightly CI?

alex-breslow-amd · 2026-09-24T12:37:22Z

There is no check that nWarps * 18KiB <= args->maxDynamicSmem. A static_assert on sizeof(tmaSmemStruct) vs 18 KiB, plus a device-side compare to maxDynamicSmem, would catch the next tile widening. The PR description still says “8 warps x 10KiB", no resize"; that is no longer true.

Some of the device-side checks were overly expensive, so I lifted them to the host. However, I was able to add a device-side static assert. Functionally in terms of correctness checking, it should be roughly equivalent.

amd-argus

ARGUS

Bottom line: Worth settling before approval: ncclSymkTmaKernelMask() at symmetric_sched.cc:349 has no fake under test/host/fakes, so rccl-UnitTestsMicroWarpSpeed stops linking. The 42-stub link break from the last round is still open too.

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

amd-argus

ARGUS

This PR changes 1,073 lines across 20 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=2 narrows kmask to the Tma bits and the GIN filter then zeroes it on any multi-node comm, dropping the symmetric path. The force path's own tests all skip when the SDK ships no TDM header.

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

mustafabar · 2026-09-24T20:55:33Z

LGTM

The symmetric staging wrappers always passed Aligned=false, so every tile issued head/bulk/tail descriptors even when both sides already sat on a 128 B line. ncclSymkTileLoad/ncclSymkTileStore gain a TileAligned parameter and AllGather's 256 B tier passes it: nPreBytes peels input to 256 B, the alignment % 256 gate carries output with it, and every tile advances by a multiple of 256, so one bulk descriptor replaces three. The scratch declaration gains __align__(128), which documents what the window placement already relied on -- buff is at offset 0 and the per-warp stride is 144 x 128 -- rather than changing any address. Keeping it off tmaSmemStruct matters: alignas(128) there moves nothing but raises alignof, and that alone cost AllReduce and ReduceScatter 7-12% through worse codegen on their per-lane tile reads, which AllGather does not perform. Measured on 4x gfx1250, TDM forced, out-of-place bus bandwidth: AllGather +23% at 16 MB and +16% at 64 MB; AllReduce and ReduceScatter within +-2%. Not a correctness fix. tdmCopy.h's head peel handles any alignment on either side, verified over 448 mismatched global/LDS combinations and down to 1-byte LDS offsets, so the review finding it answers does not hold. Still open: AllGather loses 7-12% at a 20-byte user offset, where the 256 B tier should be gated off entirely and this change should be inert. Co-authored-by: Cursor <cursoragent@cursor.com>

…hrowing agents off.

amd-argus

ARGUS

This PR changes 1,126 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: Nothing new worth raising; still open from earlier: test_symmetric_kernels.py:64 and 5 more. Worth clearing before approval.

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

warpLdsCopy()'s bulk used a 2-D descriptor: TD0 elements across, `rows` down, with tensorDim0Stride == TD0. The rows are contiguous, so that stride makes the second dimension describe exactly the same bytes as a flat run of rows*TD0 elements. The extra dimension bought nothing. ldsTileRows() now emits that flat run, and both it and ldsTileBytes() leave tileDim1, tensorDim1 and tensorDim0Stride at their zero-initialised values -- a one-row tile has no second dimension to describe and no row pitch to stride by. The descriptor zero-initialises and the setters store values verbatim with no N-1 encoding, so these really are left as 0 rather than defaulting to 1; the hardware accepts that. tile_dim0 now carries the whole length instead of one row's worth, so a single bulk descriptor tops out at MAX_ROWS = TILE_DIM_MAX / TD0 = 1023 rows (256 KB) where it was 65535 rows before. The bulk loop chunks against that. Nothing in the symmetric kernels comes near it -- the widest tile they stage is 16 KB. Scope is the one-way global<->LDS path only. issueRows()/issueRow1d(), which serve the round-trip tdmCopy*() entry points on the SIMPLE protocol, keep their existing shape. Micro-benchmark, one warp, LDS side held aligned so only the global residue varies, median of 3: at 32 KB the dword-aligned-but-not-16 offsets lose their irregular penalty -- offsets 4, 12 and 124 go from 1.69/1.71/1.83 us to 1.13/1.13/1.25, about a third faster -- and 16/32/128 B offsets gain ~8%. Aligned costs 1%, and sub-dword offsets cost ~9% on a path already 16x off aligned in both shapes, which is the DS4 data size rather than the tile form. End-to-end on the symmetric kernels with TDM forced (4x gfx1250, out-of-place bus bandwidth, median of 3, -u 4 = 16 B user offset): ReduceScatter 64M +36.5% 256M +54.7% 1G +64.4% (774 -> 1273 GB/s at 1G) AllReduce 64M +20.9% 256M +38.5% 1G +56.0% (760 -> 1186 GB/s at 1G) At a 20 B offset the same pattern holds, +52% for both at 1 GB. This closes most of the pre-existing misalignment cliff: ReduceScatter at 1 GB now reaches 1273 GB/s against 1325 on aligned buffers, where before it managed 774. Aligned buffers are unaffected -- AllReduce and ReduceScatter within +-1.8% at every size, AllGather positive above 16 MB. The 16 MB row carries 8-30% run-to-run spread in this harness and is not load-bearing for any claim above; 64 MB and up sit at <=2%. Co-authored-by: Cursor <cursoragent@cursor.com>

systems-assistant · 2026-09-28T04:39:44Z

TheRock Submodule Bump Activity

Newest first

2026-09-28 — Included in TheRock via ROCm/TheRock#8533.

## 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 1. Host microtests: build and run `rccl-UnitTestsMicro`, `rccl-UnitTestsMicroSymKernels`, `rccl-UnitTestsMicroTaskPrep`, `rccl-UnitTestsMicroEnqueue`, `rccl-UnitTestsMicroInit` and `rccl-UnitTestsMicroCe`. 2. Host microtests for the in-place alignment rule: `SchedulerMicrotest.SymBatchAligned16B_InPlace*` in `rccl-UnitTestsMicroWarpSpeed`. 3. 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. 4. 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 - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Alex Breslow <abreslow@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Praneeth Vallem <pvallem@amd.com>

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

therock-pr-bot Bot added the Not ready to Review label Sep 18, 2026

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

alex-breslow-amd force-pushed the users/abreslow/tmaSymmetricPort branch from 61fc2fd to 87dd3b2 Compare September 18, 2026 21:32

alex-breslow-amd changed the base branch from develop to users/abreslow/tdmOneWay September 19, 2026 14:57

alex-breslow-amd added this pull request to stack #11900 September 19, 2026 14:57

alex-breslow-amd force-pushed the users/abreslow/tdmOneWay branch from a5e20f6 to 707fbeb Compare September 21, 2026 13:28

alex-breslow-amd force-pushed the users/abreslow/tmaSymmetricPort branch from 08f7b27 to 856feef Compare September 21, 2026 13:28

alex-breslow-amd removed this pull request from stack #11900 September 21, 2026 17:37

alex-breslow-amd changed the base branch from users/abreslow/tdmOneWay to develop September 21, 2026 17:38

alex-breslow-amd force-pushed the users/abreslow/tmaSymmetricPort branch from c1cab46 to 77696df Compare September 21, 2026 17:49

alex-breslow-amd marked this pull request as ready for review September 21, 2026 18:00

alex-breslow-amd requested review from a team and a lite review from Copilot September 21, 2026 18:00

Copilot started reviewing on behalf of alex-breslow-amd September 21, 2026 18:01

View session

Copilot AI reviewed Sep 21, 2026

View reviewed changes

Comment thread projects/rccl/src/device/symmetric/primitives.cuh Resolved

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

Comment thread projects/rccl/src/tuning/sym_model.cc Resolved

amd-argus Bot reviewed Sep 21, 2026

View reviewed changes

alex-breslow-amd commented Sep 22, 2026

View reviewed changes

Comment thread projects/rccl/src/device/symmetric/all_gather.cuh Resolved

mustafabar reviewed Sep 22, 2026

View reviewed changes

Comment thread projects/rccl/src/device/symmetric/generate.py Resolved

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

Comment thread projects/rccl/src/device/symmetric/all_gather.cuh Resolved

charest requested review from garzaran and nusislam September 23, 2026 15:56

amd-argus Bot reviewed Sep 23, 2026

View reviewed changes

alex-breslow-amd and others added 2 commits September 23, 2026 22:50

Use device scope where feasible instead of system scope

0e908e5

alex-breslow-amd and others added 4 commits September 23, 2026 22:52

Add tuning model changes

7f01960

Improve tuning logic and performance. ~+35% uptick for ReduceScatter

a13ac61

Remove overly verbose comment

fc830ce

alex-breslow-amd force-pushed the users/abreslow/tmaSymmetricPort branch 2 times, most recently from a81c1e9 to fe74859 Compare September 23, 2026 23:14

amd-argus Bot reviewed Sep 24, 2026

View reviewed changes

sdarisala approved these changes Sep 24, 2026

View reviewed changes

sdarisala reviewed Sep 24, 2026

View reviewed changes

Address Nusrat's feedback

a7988c2

amd-argus Bot reviewed Sep 24, 2026

View reviewed changes

Comment thread projects/rccl/src/scheduler/symmetric_sched.cc Resolved

Add more tests. Address feedback

6bf6555

amd-argus Bot reviewed Sep 24, 2026

View reviewed changes

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

mustafabar self-requested a review September 24, 2026 20:55

mustafabar approved these changes Sep 24, 2026

View reviewed changes

alex-breslow-amd and others added 2 commits September 25, 2026 14:33

Fix wording in documentation in header for TDM library because it's t…

f54d703

…hrowing agents off.

amd-argus Bot reviewed Sep 25, 2026

View reviewed changes

alex-breslow-amd and others added 2 commits September 25, 2026 15:42

Merge branch 'develop' into users/abreslow/tmaSymmetricPort

97c6cc5

alex-breslow-amd merged commit af6d8a1 into develop Sep 26, 2026
85 of 90 checks passed

alex-breslow-amd deleted the users/abreslow/tmaSymmetricPort branch September 26, 2026 01:53

kuhar mentioned this pull request Sep 29, 2026

[rocjitsu] Fix floating-point and operand semantics #12467

Merged

alex-breslow-amd mentioned this pull request Oct 13, 2026

feat(rccl): enable gfx1250 TDM symmetric kernels by default #13357

Merged

1 task done

Conversation

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

Uh oh!

Overview

Some Details

Issue Tracking

Test Plan

Test Result

Not covered

Submission Checklist

Uh oh!

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

Uh oh!

✅ All Policy Checks Passed

Uh oh!

therock-pr-bot Bot commented Sep 18, 2026

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!

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!

mustafabar left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

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!

nusislam commented Sep 23, 2026

Uh oh!

nusislam commented Sep 23, 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!

sdarisala left a comment

Choose a reason for hiding this comment

Uh oh!

alex-breslow-amd commented Sep 24, 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!

amd-argus Bot left a comment

alex-breslow-amd commented Sep 18, 2026 •

edited

Loading

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

edited

Loading

mustafabar left a comment •

edited

Loading

alex-breslow-amd commented Sep 24, 2026 •

edited

Loading