alex-breslow-amd · Oct 23, 2026 · Oct 13, 2026 · Oct 13, 2026 · Oct 13, 2026 · Oct 13, 2026
@@ -53,6 +53,8 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
   * DDA VMM `ncclAllReduce` and `ncclReduceScatter` launch 2048 / nRanks blocks, clamped to 32 to 128 and never more than before.
 * `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.
 * 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.
+* `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.
+* 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.
 
 ### Removed
 * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it.
@@ -102,7 +104,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Multiple GIN proxy progress threads via `NCCL_GIN_PROXY_NTHREADS`.
 * Enabled hierarchical Copy Engine `ncclAllGather`, `ncclAlltoAll`, and `ncclAlltoAllv` on multi-node communicators: the intra-node phase runs over SDMA/CE across the LSA team and the inter-node phase over one-sided RMA through the CPU (GIN) proxy. The path is off by default and takes two opt-ins: `NCCL_CTA_POLICY=2` (equivalently `config.CTAPolicy = NCCL_CTA_POLICY_ZERO`), and symmetric registration of both the send and the receive buffer through `ncclCommWindowRegister`. On architectures other than gfx1250 (MI450) a third is needed, `NCCL_CUMEM_ENABLE=1`, because symmetric memory depends on cuMem and cuMem auto-detection enables itself only on gfx1250. `NCCL_GIN_ENABLE`, `NCCL_NUM_RMA_CTX`, `NCCL_WIN_ENABLE` and `NCCL_DMABUF_ENABLE` need no change, since their defaults already qualify; the inter-node rail draws from the internal RMA context pool (`NCCL_NUM_RMA_INT_CTX`, default 4) rather than the user-addressable contexts, so `numRmaCtx` only has to stay above zero. Selection also requires an available GIN backend and the same number of ranks on every node; communicators that do not qualify keep using the kernel path. `ncclAlltoAllv` walks the gathered per-peer size matrix, skips zero-length peers, and chunks each inter-node put; single-node AlltoAllv stays on the existing CE path. Validated for correctness and path selection from 1 to 32 nodes.
 * Added `ncclGinFenceLevel` barrier semantics on the Anvil SDMA GIN backend (`NCCL_GIN_TYPE=7`). `gin.get`, `flushAsync` and `wait` are implemented instead of trapping, and a barrier's signal is now ordered behind the payload on the peer's SDMA queue, so `Put`, `Get` and the default `Put | Get` fences drain what they promise. `Put` covers a rank's puts to itself, and `ncclGinAllContexts` fences every GIN context. Single-node validated on gfx950.
-* Added gfx1250 (MI450) Tensor Data Mover staging to the symmetric AllGather, AllReduce and ReduceScatter kernels, where NCCL uses NVIDIA's TMA. Their deep loops move each tile through the warp's shared-memory scratch with the DMA engine instead of per-lane vector loads; on ROCm that path previously compiled out entirely, so the `Tma` kernels were never generated. Launch tuning was also corrected so these kernels can use up to a quarter of a GPU's WGPs, where only 14 were reached before, which is where most of the large-message gain comes from. Opt-in via `NCCL_SYM_TMA_ENABLE=1`, and reached only when the application opts into symmetric memory. `NCCL_SYM_TMA_ENABLE=2` additionally forces these kernels wherever the collective has one, skipping both the tuner's cost comparison and the deep-loop size bar, which is intended for A/B measurement rather than production. The new kernels are generated only when gfx1250 is in `GPU_TARGETS`, and every other architecture keeps the vector path unchanged.
+* Added gfx1250 (MI450) Tensor Data Mover staging to the symmetric AllGather, AllReduce and ReduceScatter kernels, where NCCL uses NVIDIA's TMA. Their deep loops move each tile through the warp's shared-memory scratch with the DMA engine instead of per-lane vector loads; on ROCm that path previously compiled out entirely, so the `Tma` kernels were never generated. Launch tuning was also corrected so these kernels can use up to a quarter of a GPU's WGPs, where only 14 were reached before, which is where most of the large-message gain comes from. Reached only when the application opts into symmetric memory; `NCCL_SYM_TMA_ENABLE` now defaults to `1` (see Changed in RCCL 2.32.3). `NCCL_SYM_TMA_ENABLE=2` additionally forces these kernels wherever the collective has one, skipping both the tuner's cost comparison and the deep-loop size bar, which is intended for A/B measurement rather than production. The new kernels are generated only when gfx1250 is in `GPU_TARGETS`, and every other architecture keeps the vector path unchanged.
 * Removed GIN rocSHMEM GDA dependency on rocSHMEM GDA bitcode: GIN rocSHMEM GDA is now header-only. The CMake `roc::rccl` target provides the required include directories in its `INSTALL_INTERFACE` property.
 
 ### Changed

@@ -1155,3 +1155,21 @@ threshold table introduced for gfx1250.
         | ``0``: No lower-bound suppression (symk eligible at all sizes).
         | ``N`` (bytes): Below ``N``, prefer DDA over the symmetric kernel.
 
+    * - | ``NCCL_SYM_TMA_ENABLE``
+        | Controls the Tensor Data Mover (TDM) symmetric kernels
+          (``AllGather_TmaST``, ``AllReduce_RSxTmaLD_AGxTmaST`` and
+          ``ReduceScatter_TmaLD``), which stage data through the gfx1250 Tensor
+          Data Mover instead of per-lane vector loads and stores. They run only on
+          gfx1250 builds whose ROCm SDK ships the TDM descriptor header, and only
+          for collectives on symmetric registered windows whose buffers are
+          16-byte aligned. On every
+          other architecture this variable has no effect. When enabled, the
+          symmetric cost model offers these kernels only at message sizes where
+          their TDM deep loop engages, and launches them with the widest allowed
+          block count. See :ref:`symmetric-tdm-kernels`.
+      - | ``0``: Disabled. The symmetric path uses only the vector-load kernels.
+        | ``1``: Offered to the symmetric cost model (default).
+        | ``2``: Forced. Wherever the collective has a TDM kernel, it is selected
+          regardless of predicted time or message size. Intended for A/B
+          measurement, not production.
+
@@ -79,7 +79,7 @@ explicitly designed for downstream forks.
 | Per-collective CTA/CGA and `CTAPolicy` override | **Adapt** | `cgaClusterSize` is a Hopper+ cluster-dimension launch attribute with no HIP equivalent; accept and ignore it rather than rejecting it, so portable applications still run |
 | `userProfilerTag` plumbed to profiler | **Port** | `69e7c2c0`, `0dab2642`, `deb1f27a`. Depends on profiler v7 below |
 | Unified cost model (`src/tuning/cost_model.{cc,h}`, `src/include/tuning.h`) | **Adapt** | Replaces `ncclTopoGetAlgoTime` plus the separate `sym_kernels.cc` `queryModel_*` path with one `ncclTuningInput_t`/`ncclTuningResult_t` interface covering legacy, symmetric and CE methods. RCCL's CSV tuner and `rccl_wrap.cc` overrides must be re-expressed against it |
-| TMA kernels enabled by default and costed | **Skip** | RCCL pins `NCCL_PARAM(SymTmaEnable, "SYM_TMA_ENABLE", 0)` in `src/sym_kernels.cc`; the cost-model hooks come along with the refactor but stay inert |
+| TMA kernels enabled by default and costed | **Adapt** | On gfx1250, symmetric registered collectives can now select the Tensor Data Mover (`Tma`) kernels by default, and the cost model picks them for large messages. `NCCL_SYM_TMA_ENABLE=0` turns them off, and `2` forces them for A/B measurement. Other architectures are unaffected |
 | Algorithm selection on newer Intel CPUs | **Defer** | Bandwidth-table tuning only. Low value on AMD platforms; revisit if Intel-host MI systems matter |
 | AMD EPYC topology modeling (`#2036`) | **Port** | Community contribution that improves AMD host modeling. RCCL currently classifies AMD CPUs only as `ZEN`/`ROME` in `src/graph/topo.cc`. Directly relevant despite arriving via upstream |
 

@@ -78,6 +78,68 @@ missing (for example, ``cuMemEnable`` or ``globalGinSupport``). If you expect
 the symmetric path but it is disabled, check those prerequisites rather than
 ``NCCL_P2P_LEVEL``.
 
+.. _symmetric-tdm-kernels:
+
+Tensor Data Mover symmetric kernels on gfx1250
+----------------------------------------------
+
+On gfx1250 (MI450), the symmetric allgather, allreduce, and reduce-scatter
+kernels have a variant that stages data through the Tensor Data Mover (TDM)
+rather than per-lane vector loads and stores. TDM is a DMA engine that moves
+tiles between global memory and the LDS (shared memory). As in NCCL 2.31.2,
+these kernels are enabled by default (``NCCL_SYM_TMA_ENABLE=1``).
+
+You don't need to change anything in your application to use them. A collective
+can use a TDM kernel when all of the following are true:
+
+- RCCL was built with ``gfx1250`` in ``GPU_TARGETS`` against a ROCm SDK that
+  ships the TDM descriptor header (``hip/amd_detail/amd_gfx1250_TDM.h``). The
+  TDM kernels are not generated otherwise.
+- The communicator can use symmetric memory (see the prerequisites above), and
+  the send and receive buffers are registered as symmetric windows, as for any
+  symmetric kernel.
+- The operation is an allgather of any datatype, or an allreduce with
+  ``ncclSum`` or a reduce-scatter with ``ncclSum`` or ``ncclAvg`` on a
+  floating-point type (``float``, ``half``, ``bfloat16``, ``fp8_e4m3``, or
+  ``fp8_e5m2``).
+- The send and receive buffers sit at the same offset, modulo 16 bytes, within
+  their windows. For an in-place allgather or reduce-scatter, the per-rank slice
+  must also be a multiple of 16 bytes.
+- The message is large enough for the TDM deep loop to engage.
+
+The symmetric cost model then chooses between the TDM kernel and the vector-load
+kernels. The TDM kernels cost more setup latency but reach a higher bandwidth
+ceiling, so the model picks them for large messages. When it does, the
+kernel launches at 512 threads (16 waves) per block and the widest block count
+allowed, because TDM throughput keeps scaling with resident blocks.
+
+To check which kernel ran, enable tuning logs and look for the ``[Symmetric]``
+line. TDM kernels have ``Tma`` in their name:
+
+.. code-block:: shell
+
+   NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=TUNING ./your_app
+   # ... ReduceScatter [Symmetric]: 67108864 Bytes -> Kernel ReduceScatter_TmaLD nchannels 64 nthreads 512 nWorks 1
+
+In rccl-tests, ``-A 1`` prints the selected algorithm for each size. With
+symmetric registration (``-R 2``), the ``algo`` column shows ``SymTDM`` for sizes
+that use a TDM kernel and ``SYM`` for sizes that use a vector-load symmetric
+kernel:
+
+.. code-block:: shell
+
+   ./reduce_scatter_perf -g 4 -b 64M -e 256M -f 4 -R 2 -A 1
+   #       size  ...   algo     proto   nchannels
+   #   67108864  ...    SYM    SIMPLE          62
+   #  268435456  ...  SymTDM    SIMPLE          64
+
+``NCCL_SYM_TMA_ENABLE`` controls these kernels:
+
+- ``0`` disables them.
+- ``1`` (the default) offers them to the cost model.
+- ``2`` forces them wherever the collective has one, ignoring the cost model and
+  the minimum size. Use ``2`` for A/B measurements only.
+
 .. _device-symmetric-memory:
 
 Device-side symmetric memory access (``ncclFindWindow``)

@@ -62,11 +62,8 @@ static ncclResult_t fillCollTuningInput(struct ncclComm* comm, struct ncclRawTas
   }
   in->nvlsSupport = ncclNvlsTransportEnabled(comm) &&
                     (ncclNvlsSupported(raw->opDev.op, raw->datatype) || raw->func == ncclFuncAllGather);
-  {
-    size_t inputOff = sendWin ? (uintptr_t)raw->sendbuff - (uintptr_t)sendWin->userPtr : (uintptr_t)raw->sendbuff;
-    size_t outputOff = recvWin ? (uintptr_t)raw->recvbuff - (uintptr_t)recvWin->userPtr : (uintptr_t)raw->recvbuff;
-    in->symAligned16B = (uint32_t(inputOff - outputOff) % 16 == 0);
-  }
+  in->symAligned16B = ncclSymAligned16B(raw->func, comm->rank, raw->sendbuff, raw->recvbuff, sendWin, recvWin,
+                                        raw->count * ncclTypeSize(raw->datatype));
 
   sendbuffSize = elementSize * ncclFuncSendCount(raw->func, comm->nRanks, raw->count);
   recvbuffSize = elementSize * ncclFuncRecvCount(raw->func, comm->nRanks, raw->count);

@@ -115,6 +115,23 @@ static inline bool ncclAllGatherIsInPlace(const void* sendbuff, const void* recv
   return (uintptr_t)sendbuff == (uintptr_t)recvbuff + size_t(rank) * perRankBytes;
 }
 
+// [RCCL] Whether the Tma symmetric kernels may run: input and output must sit at the same offset
+// modulo 16 within their windows. In-place AllGather and ReduceScatter shift one side by
+// rank * perRankBytes, so that check alone can differ across ranks, and ranks that pick different
+// symmetric kernels launch grids that never meet at the LSA barrier. There the per-rank slice must
+// itself be 16 B aligned, so every rank reaches the same answer. Null buffers count as aligned.
+static inline bool ncclSymAligned16B(ncclFunc_t func, int rank, const void* sendbuff, const void* recvbuff,
+                                     struct ncclDevrWindow* sendWin, struct ncclDevrWindow* recvWin,
+                                     size_t perRankBytes) {
+  size_t inputOff = sendWin ? (uintptr_t)sendbuff - (uintptr_t)sendWin->userPtr : (uintptr_t)sendbuff;
+  size_t outputOff = recvWin ? (uintptr_t)recvbuff - (uintptr_t)recvWin->userPtr : (uintptr_t)recvbuff;
+  if (uint32_t(inputOff - outputOff) % 16 != 0) return false;
+  if (sendbuff == nullptr || recvbuff == nullptr) return true;
+  bool inPlace = (func == ncclFuncAllGather && ncclAllGatherIsInPlace(sendbuff, recvbuff, rank, perRankBytes)) ||
+                 (func == ncclFuncReduceScatter && ncclAllGatherIsInPlace(recvbuff, sendbuff, rank, perRankBytes));
+  return !inPlace || perRankBytes % 16 == 0;
+}
+
 ncclResult_t ncclGetCollNetSupport(struct ncclComm* comm, struct ncclTaskColl* task, int* collNetSupport);
 ncclResult_t ncclGetAlgoInfo(struct ncclComm* comm, struct ncclTaskColl* task, int collNetSupport, int nvlsSupport,
                              int numPipeOps, ncclSimInfo_t* simInfo = NULL);

@@ -89,6 +89,9 @@ typedef enum {
   // Appended rather than grouped with the other Direct entries so existing
   // values stay stable.
   RCCL_DIRECT_ALLTOALL, // AlltoAll as per-peer Send/Recv (no collective kernel)
+  // Reporting only (rcclGetCollImplInfo / rcclSymKGetInfo): a symmetric dispatch that
+  // resolves to a Tensor Data Mover kernel. Live dispatch always tags RCCL_SYMMETRIC.
+  RCCL_SYMMETRIC_TDM,
   RCCL_ALGO_COUNT
 } rcclAddonAlgos_t;
 

@@ -221,6 +221,7 @@ const char* ncclSymkKernelIdToString(int kernelId);
 ncclResult_t ncclGetSymRegType(struct ncclDevrWindow* sendWin, struct ncclDevrWindow* recvWin,
                                ncclSymRegType_t* winRegType);
 bool rcclSymkKernelIdIsLL(int kernelId);
+bool rcclSymkKernelIdIsTma(int kernelId);
 
 int ncclSymkLLKernelMask();
 int ncclSymkDynamicSmemKernelMask();