ScottTodd · Sep 28, 2026 · Sep 22, 2026 · Sep 24, 2026 · Sep 24, 2026
@@ -579,15 +579,30 @@ testResult_t AlltoAllDeviceTime(struct threadArgs* args, ncclDataType_t type, nc
 }
 #endif
 
+testResult_t AlltoAllGetAlgoProtoChannels(ncclComm_t comm, size_t count, ncclDataType_t type, int* algo, int* proto, int* nchannels) {
+  if (rcclTestsGetAlgoInfo == NULL) return testInternalError;
+  NCCLCHECK(rcclTestsGetAlgoInfo(comm, ncclFuncAlltoAll, count, type, 0, 0, 1, algo, proto, nchannels));
+  return testSuccess;
+}
+
+testResult_t AlltoAllGetCollImplInfo(ncclComm_t comm, size_t count, ncclDataType_t type, ncclRedOp_t op,
+    const void* sendbuff, void* recvbuff, int graphCapturing, int* algo, int* proto, int* nchannels) {
+  if (rcclTestsGetCollImplInfo == NULL) return testInternalError;
+  NCCLCHECK(rcclTestsGetCollImplInfo(comm, ncclFuncAlltoAll, count, type, op, sendbuff, recvbuff, graphCapturing, algo, proto, nchannels));
+  return testSuccess;
+}
+
 struct testColl alltoAllTest = {
   "AlltoAll",
   AlltoAllGetCollByteCount,
   AlltoAllInitData,
   AlltoAllGetBw,
   AlltoAllRunColl,
+  AlltoAllGetAlgoProtoChannels,
+  // No symk hook: the symmetric kernels cover AllReduce/AllGather/ReduceScatter
+  // only, so rcclSymKGetInfo can never describe an AllToAll.
   NULL,
-  NULL,
-  NULL,
+  AlltoAllGetCollImplInfo,
   AlltoAllDeviceTime
 };
 

@@ -454,18 +454,30 @@ testResult_t testLaunchDeviceKernel(F kernel, void* sendbuff, size_t sendoffset,
 #define SPECIALIZE_KERNEL(kernel, type, op) nullptr
 #endif
 
+// Must match librccl's internal ncclFunc_t (src/include/nccl_common.h): these
+// values cross the rcclGetAlgoInfo / rcclSymKGetInfo / rcclGetCollImplInfo ABI,
+// so a renumbering here silently mislabels results instead of failing to build.
 typedef enum {
   ncclFuncBroadcast = 0,
   ncclFuncReduce = 1,
   ncclFuncAllGather = 2,
   ncclFuncReduceScatter = 3,
   ncclFuncAllReduce = 4,
-  ncclFuncAllReduceWithBias = 5,
-  ncclFuncSendRecv = 6,
-  ncclFuncSend = 7,
-  ncclFuncRecv = 8,
-  ncclFuncAllToAllPivot = 9,
-  ncclNumFuncs = 10
+  ncclFuncSendRecv = 5,
+  ncclFuncSend = 6,
+  ncclFuncRecv = 7,
+  ncclFuncAlltoAll = 8,
+  ncclFuncScatter = 9,
+  ncclFuncGather = 10,
+  ncclFuncAllToAllPivot = 11,
+  ncclFuncAlltoAllGda = 12,
+  ncclFuncAlltoAllvGda = 13,
+  ncclFuncAllGatherV = 14,
+  ncclFuncPutSignal = 15,
+  ncclFuncSignal = 16,
+  ncclFuncWaitSignal = 17,
+  ncclFuncAlltoAllv = 18,
+  ncclNumFuncs = 19
 } ncclFunc_t;
 
 typedef ncclResult_t (*rcclTestsGetAlgoInfo_t)(struct ncclComm* comm, ncclFunc_t coll, uint64_t count, ncclDataType_t dataType,

@@ -5,6 +5,13 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 ## RCCL 2.31.2 for ROCm 10.0.0 (Unreleased)
 
 ### Added
+* `RCCL_CE_AR_MAX_MSG_BYTES` (default `-1`): overrides the 2-shot AllReduce size cap from the arch table. Set to a positive value to override `ceNonRegMax[AR]`.
+* `RCCL_CE_AR_REG_MAX_MSG_BYTES` (default `-1`): overrides the registered CE AllReduce size cap from the arch table.
+* `RCCL_CE_AR_STAGING_BYTES` (default `-1`): overrides the CE AllReduce staging buffer allocation size; when unset, `NCCL_CE_AR_STAGING_BYTES` is used.
+* Per-architecture dispatch table (`rcclArchThresholds`) centralizing algo/proto selection for DDA protocol, CE, and symmetric kernel per collective, replacing scattered hardcoded defaults. Tables are defined for gfx1250, gfx950, and gfx942; the lookup `rcclGetArchThresholds()` maps GCN arch strings to the appropriate entry. The `symMinR2[func]` field sets a per-collective lower bound for registered buffers.
+* `RCCL_IGNORE_ARCH_TABLE` (default `0`): when set to `1`, bypasses the `rcclArchThresholds` dispatch table and falls back to compile-time constants for all DDA, CE, and symmetric-kernel thresholds. Has no effect on gfx942 and gfx950, which always use the arch table.
+* New `rcclAddonAlgos_t` values: `RCCL_CE_SCRATCH` (CE via DDA scratch, distinct from `RCCL_CE_REGISTERED`), `RCCL_A2A_PIVOT`, `RCCL_A2A_GDA`, `RCCL_A2A_GIN_SDMA`, and `RCCL_DIRECT_ALLTOALL`, with corresponding `rcclGetAlgoName()` labels.
+* rccl-tests: AlltoAll now reports algo/protocol in VERSION output, consistent with other collectives.
 * Compatibility with NCCL 2.31.2.
 * Per-collective configuration APIs (`ncclCollConfig_t` / `nccl*Config()` entry points) and the `ncclConfigExt_t` vendor extension list. Initialize configs with `NCCL_COLLCONFIG_INITIALIZER`.
 * Communicator config (`ncclConfig_v23100`) fields for implicit launch ordering (`launchOrderImplicit`), RMA signal count (`numRmaSig`), eager RMA init (`rmaEagerInit`), and host collective fault tolerance (`hostCftMode`).
@@ -15,6 +22,21 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Enabled hierarchical Copy Engine `ncclAllGather` and `ncclAlltoAll` 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. Validated for correctness and path selection from 1 to 32 nodes.
 
 ### Changed
+
+* `rcclSelectAllGather` signature changed — it now takes a `cudaStream_t stream` parameter (for live dispatch) and a `bool query` mode that uses `graphCapturingHint` instead of probing the stream. The comment block was updated too.
+* `RCCL_DDA_THRESHOLD`, `RCCL_DDA_LL_THRESHOLD`, and `RCCL_DDA_LL128_THRESHOLD` now default to `-1` (unset), resolved at runtime from the arch dispatch table instead of being hardcoded to 128 MiB and 64 KiB respectively (RCCL_DDA_LL128_THRESHOLD had no pre-table default; the LL128 tier was off by default). Setting these env vars explicitly still takes precedence over the table.
+* `RCCL_DDA_LL128` now defaults to `-1` (auto): the LL128 tier is enabled only for architectures whose arch-table row has a non-zero `ddaLL128Max` for the collective (currently gfx1250 only).
+* `RCCL_CE_ALLREDUCE` now defaults to `-1` (auto/enabled) instead of `0` (disabled). CE AllReduce is therefore on by default for gfx1250 communicators that meet all other eligibility criteria.
+* CE-2-Shot AllReduce size cap is now resolved at runtime via `rcclCeAr2ShotMax()` (env var wins, then arch table `ceNonRegMax[AR]`, then the 256 MiB `NCCL_CE_AR_TMPBUF_DEFAULT_BYTES` fallback) rather than from a compile-time constant.
+* `RCCL_FORCE_CE_ALLREDUCE=1` no longer overrides the CE AllReduce 2-shot selector cap. It bypasses the `CTA_POLICY_ZERO` check only; the selector cap is enforced regardless.
+* `tuning_model_11` registered for gfx1250 single-node; gfx1250 added to the AINIC tuning index map.
+* Extended `ncclFunc_t` enum in rccl-tests `common.h` to match the librccl internal ABI: renumbered existing values and added `ncclFuncAlltoAll` (8), `ncclFuncScatter` (9), `ncclFuncGather` (10), `ncclFuncAllToAllPivot` (11), `ncclFuncAlltoAllGda` (12), `ncclFuncAlltoAllvGda` (13), `ncclFuncAllGatherV` (14), `ncclFuncPutSignal` (15), `ncclFuncSignal` (16), `ncclFuncWaitSignal` (17), `ncclFuncAlltoAllv` (18); `ncclNumFuncs` updated to 19. 
+* AlltoAll index variables in `alltoall_dda.h` and `alltoall_dda_fabric.h` widened from `int` to `size_t` to prevent overflow on large message sizes.
+* DDA fabric and IPC AlltoAll launchers refactored to derive launch geometry from a single inline helper (`ddaAllToAllFabricGeom`, `ddaAllToAllFabricLLGeom`, `ddaAllToAllFabricLL128Geom`, `ddaAllToAllIpcGeom`) shared between the dispatch path and the new block-count query functions.
+* `rcclGetCollImplInfo` now covers ReduceScatter and AlltoAll (previously both fell back to `rcclGetAlgoInfo()`); rccl-tests reports accurate algo/proto for all four collective types.
+* Algorithm selectors (`rcclSelectAllReduce`, `rcclSelectAllGather`, `rcclSelectReduceScatter`, `rcclSelectAlltoAll`) now tag the task with `rcclSymkExtract` (`ALLOW`/`DENY`) so `taskAppend` does not re-derive the symmetric kernel decision. Setting `NCCL_ALGO=Ring` now correctly prevents the symmetric kernel from being extracted even when both buffers are registered.
+
+
 * **Breaking: `NCCL_GIN_TYPE` values for AMD backends are not compatible with 2.30.7.** NCCL 2.31 inserted EFA GDA at value 5, so rocSHMEM GDA moved 5→6 and Anvil SDMA moved 6→7. The IB proxy remains `2`. Jobs that still set `NCCL_GIN_TYPE=6` now select rocSHMEM GDA, not Anvil SDMA. See `src/gin/README.md`.
 * One-sided RMA supports multiple contexts and signals; the previous restriction to context 0 and signal index 0 has been lifted (`numRmaCtx` / `numRmaSig`).
 * Updated the RMA plugin interface to v15.

@@ -607,3 +607,118 @@ variables are collected in the following table.
         | disabled, events are dropped once a pool is exhausted.
       - | ``0``: Fixed-size pools.
         | ``1``: Pools grow on demand (default).
+
+Algorithm dispatch and tuning (gfx1250 / MI450)
+================================================
+
+The following variables control per-architecture algorithm and protocol dispatch,
+including the DDA fabric tiers, Copy Engine (CE) AllReduce, and the per-architecture
+threshold table introduced for gfx1250.
+
+.. list-table::
+    :header-rows: 1
+    :widths: 40,60
+
+    * - **Environment variable**
+      - **Values**
+
+    * - | ``RCCL_IGNORE_ARCH_TABLE``
+        | Controls whether the per-architecture dispatch table
+          (``rcclArchThresholds``) is used to set DDA, CE, and symmetric-kernel
+          thresholds, or whether pre-table compile-time constants are used instead.
+          When the table is active (``0``), each threshold variable below
+          (``RCCL_DDA_THRESHOLD``, ``RCCL_DDA_LL_THRESHOLD``, etc.) defaults to
+          ``-1`` (unset) and is resolved from the table at runtime; setting any of
+          those variables explicitly still overrides the table.
+      - | ``0``: Use the per-architecture table for threshold defaults (default).
+        | ``1``: Ignore the table; fall back to pre-table compile-time constants.
+          This matches the behavior before the arch table was introduced.
+
+    * - | ``RCCL_DDA_THRESHOLD``
+        | Upper bound in bytes for the DDA VMM/Simple tier per collective.
+          Messages above this size exit DDA and fall through to Ring, CE, or the
+          symmetric kernel. When set to ``-1`` (default), the value is resolved
+          from the per-architecture table (``ddaVmmMax`` field). Set
+          ``RCCL_IGNORE_ARCH_TABLE=0`` for the table to take effect.
+      - | ``-1``: Resolved from the per-arch table at runtime (default).
+        | ``0``: Disable the DDA VMM tier for all collectives.
+        | ``N`` (bytes): Use ``N`` as the VMM tier ceiling for all collectives,
+          overriding the table. Pre-table default: ``134217728`` (128 MiB).
+
+    * - | ``RCCL_DDA_LL_THRESHOLD``
+        | Upper bound in bytes for the DDA LL (low-latency) tier. Messages at or
+          below this size use the LL protocol; larger messages move to LL128 or VMM.
+          When ``-1`` (default), resolved from ``ddaLLMax`` in the arch table.
+      - | ``-1``: Resolved from the per-arch table at runtime (default).
+        | ``0``: Disable the DDA LL tier.
+        | ``N`` (bytes): Use ``N`` as the LL tier ceiling, overriding the table.
+          Pre-table default: ``65536`` (64 KiB).
+
+    * - | ``RCCL_DDA_LL128``
+        | Enables the DDA LL128 protocol tier. When ``-1`` (auto), LL128 is enabled
+          only for architectures whose per-arch table has a non-zero ``ddaLL128Max``
+          entry (currently gfx1250 only). ``1`` forces the tier on but still reads
+          its ceiling from the arch table, so it yields the same effective cap as
+          ``-1`` in every current configuration; ``0`` disables the tier on all
+          architectures.
+      - | ``-1``: Auto-enabled only when the arch table has non-zero LL128 thresholds
+          (default).
+        | ``0``: Disabled on all architectures.
+        | ``1``: Force-enabled; ceiling is still read from the arch table (same cap
+          as ``-1`` in every current configuration).
+    * - | ``RCCL_DDA_LL128_THRESHOLD``
+        | Upper bound in bytes for the DDA LL128 tier. Messages above this size move
+          to VMM/Simple or Ring. When ``-1`` (default), resolved from ``ddaLL128Max``
+          in the arch table. When set explicitly, it takes precedence over ``RCCL_DDA_LL128=0`` (the env var is read before the flag check).
+      - | ``-1``: Resolved from the per-arch table at runtime (default).
+        | ``0``: Disable the DDA LL128 tier.
+        | ``N`` (bytes): Use ``N`` as the LL128 tier ceiling, overriding the table.
+          Pre-table default: ``0`` (LL128 tier was off before the arch table).
+
+    * - | ``RCCL_CE_ALLREDUCE``
+        | Enables the Copy Engine (CE) registered-window AllReduce path. When ``-1`` (auto),
+          CE AllReduce is on by default for gfx1250 communicators that meet all
+          eligibility criteria, and off for all other architectures.
+      - | ``-1``: Auto — enabled on gfx1250, disabled elsewhere (default).
+        | ``0``: Disabled on all architectures.
+        | ``1``: Force-enabled (subject to other eligibility checks such as buffer
+          registration and message size).
+
+    * - | ``RCCL_FORCE_CE_ALLREDUCE``
+        | Bypasses the ``NCCL_CTA_POLICY=2`` (``CTA_POLICY_ZERO``) requirement for
+          CE AllReduce, allowing CE to run without symmetric window registration.
+          Does not override the staging buffer size cap (``RCCL_CE_AR_MAX_MSG_BYTES``
+          or the arch table ``ceNonRegMax[AR]``); that cap is enforced regardless.
+          When ``-1`` (auto), follows the same gfx1250-default logic as
+          ``RCCL_CE_ALLREDUCE``.
+      - | ``-1``: Auto — same default as ``RCCL_CE_ALLREDUCE`` (default).
+        | ``0``: Disabled.
+        | ``1``: Force-enabled (CTA_POLICY check bypassed).
+
+    * - | ``RCCL_CE_AR_MAX_MSG_BYTES``
+        | Overrides the CE 2-shot AllReduce message size cap. When ``-1`` (default),
+          the cap is read from ``ceNonRegMax[AllReduce]`` in the per-arch table. A
+          null or unknown-arch table restores the pre-table 256 MiB default. This
+          variable sizes the selector only; it does not affect the staging buffer
+          allocation (see ``RCCL_CE_AR_STAGING_BYTES``).
+      - | ``-1``: Resolved from the per-arch table (default).
+        | ``N`` (bytes): Use ``N`` as the 2-shot AllReduce size cap.
+
+    * - | ``RCCL_CE_AR_REG_MAX_MSG_BYTES``
+        | Overrides the CE registered-window AllReduce message size cap. When ``-1``
+          (default), the cap is read from ``ceRegMax[AllReduce]`` in the per-arch
+          table. This is the upper bound for the path where both send and receive
+          buffers are in symmetric windows (``-R 2`` mode in rccl-tests).
+      - | ``-1``: Resolved from the per-arch table (default).
+        | ``N`` (bytes): Use ``N`` as the registered AllReduce size cap.
+
+    * - | ``RCCL_CE_AR_STAGING_BYTES``
+        | Overrides the total allocation size of the CE AllReduce staging buffer
+          (``ceARTmpBuf``). When ``-1`` (default), the buffer is allocated at the
+          compile-time constant ``NCCL_CE_AR_STAGING_BYTES`` (16 MiB). Increasing
+          this reduces pipelining overhead for large messages but raises per-rank
+          GPU memory usage. This variable sizes the buffer only; the selector cap
+          is controlled separately by ``RCCL_CE_AR_MAX_MSG_BYTES``.
+      - | ``-1``: Use the compile-time default of 16 MiB (default).
+        | ``N`` (bytes): Set the per-slot payload capacity to ``N``; ``ceARTmpBuf`` is ``NCCL_CE_NUM_SLOTS`` (2) times that.
+
@@ -397,6 +397,7 @@ set(SRC_FILES
   include/proxy.h
   include/proxy_diag_counters.h
   include/ras.h
+  include/rccl_arch_thresholds.h
   include/rccl_common.h
   include/rccl_decision.h
   include/rccl_float8.h
@@ -522,6 +523,7 @@ set(SRC_FILES
   ras/ras_param.cc
   ras/ras_param.h
   ras/rasnet.cc
+  rccl_arch_thresholds.cc
   rccl_wrap.cc
   register/coll_reg.cc
   register/register.cc

@@ -10,6 +10,7 @@
 #include "algorithms/dda/all_gather/all_gather_dda_fabric_ll.h"
 #include "checks.h"
 #include "comm.h"
+#include "rccl_common.h"
 #include "algorithms/dda/dda_init_detail.h" // nccl_dda_detail::kDdaLLAgMaxBlocksPerPeer
 #include "debug.h"
 #include "algorithms/dda/fabric/fabric_gpu_barrier.h" // dda::common::kDdaMaxNranks
@@ -131,7 +132,7 @@ bool ncclAllGatherDdaFabricLLEligible(ncclComm* comm, const void* sendbuff, void
     return false;
   }
 
-  if (perRankBytes * (size_t)comm->nRanks > (size_t)rcclParamDdaLLThreshold()) {
+  if (perRankBytes * (size_t)comm->nRanks > rcclDdaLLThreshold(comm, ncclFuncAllGather)) {
     return false;
   }
   // One packet carries 8B of payload, so this rank's contribution has to fit the

@@ -12,7 +12,7 @@
 #include "comm.h"
 #include "debug.h"
 #include "algorithms/dda/fabric/fabric_gpu_barrier.h" // dda::common::kDdaMaxNranks
-#include "param.h"
+#include "rccl_common.h"
 
 #include <cuda_runtime.h>
 
@@ -160,7 +160,7 @@ bool ncclAllGatherDdaFabricLL128Eligible(ncclComm* comm, const void* sendbuff, v
   if ((reinterpret_cast<uintptr_t>(sendbuff) % 16) != 0 || (reinterpret_cast<uintptr_t>(recvbuff) % 16) != 0) {
     return false;
   }
-  if (perRankBytes * (size_t)comm->nRanks > (size_t)rcclParamDdaLL128Threshold()) {
+  if (perRankBytes * (size_t)comm->nRanks > rcclDdaLL128Threshold(comm, ncclFuncAllGather)) {
     return false;
   }
   // Derived from the scratch allocation

@@ -16,6 +16,7 @@
 #include "algorithms/dda/dda_init_detail.h"
 #include "algorithms/dda/fabric/fabric_gpu_barrier.h" // dda::common::kDdaMaxNranks
 #include "param.h"
+#include "rccl_common.h"
 
 #include <cuda_runtime.h>
 
@@ -157,9 +158,10 @@ bool ncclAllReduceDdaFabricLL128Eligible_old(ncclComm* comm, const void* sendbuf
   if (bytes % 8 != 0) {
     return false;
   }
-  // Use the runtime LL128 threshold (RCCL_DDA_LL128_THRESHOLD) as the cap.
-  const int64_t ll128Thresh = rcclParamDdaLL128Threshold();
-  if (ll128Thresh <= 0 || bytes > (size_t)ll128Thresh) {
+  // Cap by the resolved LL128 threshold (RCCL_DDA_LL128_THRESHOLD, else this
+  // arch's table, else the built-in default).
+  const size_t ll128Thresh = rcclDdaLL128Threshold(comm, ncclFuncAllReduce);
+  if (ll128Thresh == 0 || bytes > ll128Thresh) {
     return false;
   }
   // Scratch is sized from the actual message (compact per-call slot stride), so

@@ -48,8 +48,8 @@ __launch_bounds__(512)
 #pragma unroll NRANKS
     for (int r = 0; r < NRANKS; ++r) {
       int srcRank = r;
-      int srcIdx = idx + selfRank * idxEnd;
-      int destIdx = idx + r * idxEnd;
+      size_t srcIdx = idx + selfRank * idxEnd;
+      size_t destIdx = idx + r * idxEnd;
       *reinterpret_cast<uint4*>(&recvbuff[destIdx]) = reinterpret_cast<const uint4*>(&ipcbuffs[srcRank][srcIdx])[0];
     }
   }

@@ -47,8 +47,8 @@ __launch_bounds__(512)
 #pragma unroll kUnroll
     for (int r = 0; r < nRanksEff; ++r) {
       int srcRank = r;
-      int srcIdx = idx + selfRank * idxEnd;
-      int destIdx = idx + r * idxEnd;
+      size_t srcIdx = idx + selfRank * idxEnd;
+      size_t destIdx = idx + r * idxEnd;
       *reinterpret_cast<uint4*>(&recvbuff[destIdx]) = reinterpret_cast<const uint4*>(&ipcbuffs[srcRank][srcIdx])[0];
     }
   }