astyrrian1 · Oct 19, 2026 · Oct 13, 2026 · Oct 13, 2026 · Oct 15, 2026 · Oct 16, 2026
@@ -22,6 +22,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * `NCCL_HIER_CE_COLL_AG_RAIL_RING_ENABLE` (default `-1`): a positive value selects a ring for the inter-node rail phase of hierarchical Copy Engine `ncclAllGather`. The default keeps the direct path.
 * `NCCL_IB_SORT_MERGE_NICS`: sorts the sub-devices of a merged IB device by plane ID. RCCL defaults it to `0` (NCCL defaults to `1`), so merged-device order and names are unchanged.
 * nccl4py: per-call collective configuration (`NCCLCollConfig`, `VendorOption`) including the launch completion event, communicator properties (`NCCLCommProperties`), and the `GIN_ONLY` window flag. On ROCm, the HIP `Event` shim provides only an event handle; `record()`, `sync()` and `query` are not implemented.
+* `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.
 
 ### Changed
 * Host-side device API declarations (`ncclDevCommCreate`, `ncclDevCommDestroy`, `ncclCommQueryProperties`, `ncclGetPeerDevicePointer`, the `*CreateRequirement` helpers, the `ncclDevCommRequirements` and `ncclCommProperties` structs and their initializers) moved to `nccl_device/host.h`. `nccl_device.h` still includes it; code that includes individual `nccl_device/*.h` headers directly must also include `nccl_device/host.h`.
@@ -33,6 +34,13 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Network devices for each GPU are now chosen by rail and plane assignment, replacing the previous start-device scattering.
 * `NCCL_MLOPART_RDMA_ENABLE` (default `0`) is kept. NCCL 2.32 removed it and treats buffers on partitioned GPUs as RDMA-capable on all non-ARM hosts; RCCL keeps network buffer registration for partitioned (CPX/DPX) GPUs opt-in.
 * Retuned the symmetric AllReduce, ReduceScatter and AllGather kernels for gfx950. The block width is now selected per collective and message size instead of a fixed 256 threads, and the work partitioning and unroll factors in the load-direct kernels and the AllGather store kernel were refitted to the 64-lane wavefront. AllGather also switches from LL to the store kernel from 4 MB instead of 16 MB. The multi-node GIN kernels and all other architectures are unchanged. To fit the wider LL blocks, the shared LL scratch buffer on gfx950 grows from 4 MiB to 8 MiB per 8-rank communicator.
+* gfx1250 no longer forces unregistered `ncclAllGather`, `ncclAlltoAll`, `ncclGather` and `ncclScatter` through the Copy Engine scratch path (CE-Scratch), because it was slower than the kernel paths it replaced. Whether CE-Scratch is forced by default now comes from the per-architecture table: gfx942 and gfx950 still force it. A communicator-level `NCCL_CTA_POLICY_ZERO` still forces it on every architecture; an explicit `RCCL_FORCE_CE=1` still forces it and `RCCL_FORCE_CE=0` still turns it off. gfx1250 `ncclAllGather` now uses DDA LL128 up to 128 MiB.
+* gfx1250 `ncclAllReduce` on single-node communicators keeps the LL128 protocol at larger sizes: Tree at every size, Ring up to nRanks x 32 MiB, and the DDA LL128 tier up to 256 MiB. The DDA LL128 two-shot ceiling now follows that tier ceiling (`RCCL_DDA_LL128_TWOSHOT_THRESHOLD` still overrides it); the DDA scratch slot lets two-shot reach 256 MiB from 32 ranks, about 240 MiB at 16 and 128 MiB at 2 to 8. On gfx1250 communicators of 2 to 8 ranks, the DDA fabric scratch buffer grows to about 546 MiB.
+* On gfx1250 single-node communicators of 8 or more ranks, the DDA tier, algorithm and protocol crossovers now scale with the rank count. Communicators below 8 ranks keep the previous thresholds.
+  * `ncclAllReduce`: DDA LL two-shot up to nRanks x 1 MiB (`RCCL_DDA_LL_TWOSHOT_THRESHOLD` still overrides it); from 32 ranks, DDA VMM between 2 MiB and nRanks x 512 KiB for unregistered buffers unless `RCCL_DDA_THRESHOLD` or `RCCL_DDA_LL_THRESHOLD` is set; Tree instead of Ring up to nRanks x 1 MiB unless `NCCL_ALGO` or `NCCL_PROTO` is set, with LL up to nRanks x 32 KiB.
+  * `ncclAllGather`: DDA LL up to nRanks x 16 KiB, DDA LL128 up to 256 MiB from 16 ranks, and Ring LL128 up to 64 MiB per rank.
+  * `ncclReduceScatter`: DDA LL up to nRanks x 256 KiB, DDA LL128 up to nRanks x 512 KiB, then DDA VMM up to nRanks x 2 MiB; Ring LL up to 512 KiB per rank and LL128 up to 8 MiB per rank (32 MiB for 1-byte types); Ring Simple on at most 128 channels between 8 and 32 MiB per rank unless `NCCL_MIN_NCHANNELS` or `NCCL_MAX_NCHANNELS` is set.
+  * DDA VMM `ncclAllReduce` and `ncclReduceScatter` launch 2048 / nRanks blocks, clamped to 32 to 128 and never more than before.
 
 ### Removed
 * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it.
@@ -45,6 +53,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Fixed tuner plugins receiving uninitialized cost-model constants.
 * Fixed GIN proxy descriptor shared-memory sizing and alignment.
 * Fixed virtual address space exhaustion when symmetric windows backed by the same physical allocation are registered repeatedly.
+* Fixed Copy Engine `ncclAlltoAllv` hanging on intra-node communicators of more than 8 ranks with uneven send sizes. Each rank decided from its own send sizes whether to add intra-batch barriers, so ranks could issue different barrier counts. The decision now uses the full send-size matrix, which every rank shares.
 * Fixed `NCCL_P2P_DISABLE` and `NCCL_P2P_LEVEL` being ignored on Intel, ARM and Zhaoxin hosts, where the CPU-based default P2P level overwrote the user setting. On such hosts `NCCL_P2P_DISABLE=1` now falls back to SHM as documented, and `NCCL_P2P_LEVEL` is honored.
 
 ### Known issues

@@ -825,6 +825,19 @@ threshold table introduced for gfx1250.
         | ``1``: Ignore the table; fall back to pre-table compile-time constants.
           This matches the behavior before the arch table was introduced.
 
+    * - | ``RCCL_TUNING_PARTITION_MODE``
+        | Selects which gfx1250 per-architecture table (SPX or DPX) is used for
+          tuning thresholds. It only changes tuning, not results. A communicator
+          uses the DPX table only when every rank selects DPX; otherwise every rank
+          uses the SPX table. Other architectures have one table for every mode.
+          Has no effect when ``RCCL_IGNORE_ARCH_TABLE=1``.
+      - | ``auto`` (default): Follow the compute-partition mode the GPU reports.
+          ``DPX`` selects the DPX table; ``SPX``, ``CPX``, or no reported mode
+          selects the SPX table. Any other value is reported once and treated
+          as ``auto``.
+        | ``SPX``: Use the SPX table.
+        | ``DPX``: Use the DPX table.
+
     * - | ``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

@@ -13,6 +13,7 @@
 #include "debug.h"
 #include "algorithms/dda/dda_init_detail.h"
 #include "algorithms/dda/fabric/fabric_gpu_barrier.h"
+#include "rccl_common.h"
 
 #include <cuda_runtime.h>
 #include <hip/hip_ext.h>
@@ -29,9 +30,9 @@ using nccl_dda_detail::DdaFabricBarrierState;
 constexpr size_t kDdaFlatTreeThresholdBytes = 1ULL << 18;
 
 // Single source of the launch geometry: grid/block for `count` elements of
-// `typeSize` bytes, capped at the init-time block count.
+// `typeSize` bytes, capped at rcclDdaFabricVmmBlocks() (at most the init-time block count).
 static inline std::pair<dim3, dim3> ddaAllReduceFabricGeom(ncclComm* comm, size_t count, int typeSize) {
-  return dda::common::getGridAndBlockDims(count, typeSize, comm->ddaFabricMaxBlocks);
+  return dda::common::getGridAndBlockDims(count, typeSize, rcclDdaFabricVmmBlocks(comm));
 }
 
 template <typename T>
@@ -57,7 +58,7 @@ static ncclResult_t ncclAllReduceDdaFabricTyped(const void* sendbuff, void* recv
          count, nRanks);
   }
 
-  // Block cap chosen at init (barrier flag buffer is sized for it).
+  // At most the block cap chosen at init (barrier flag buffer is sized for it).
   auto gridBlock = ddaAllReduceFabricGeom(comm, count, sizeof(T));
   const auto& grid = gridBlock.first;
   const auto& block = gridBlock.second;

@@ -26,6 +26,7 @@
 #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>
 #include <hip/hip_ext.h>
@@ -37,9 +38,7 @@
 
 RCCL_PARAM_DECLARE(DdaLL);
 RCCL_PARAM_DECLARE(DdaLLOneShotThreshold);
-RCCL_PARAM_DECLARE(DdaLLTwoShotThreshold);
 RCCL_PARAM_DECLARE(DdaLL128OneShotThreshold);
-RCCL_PARAM_DECLARE(DdaLL128TwoShotThreshold);
 
 namespace {
 
@@ -366,7 +365,7 @@ bool ddaLLArTwoShotEligible(ncclComm* comm, const void* sendbuff, void* recvbuff
       return false;
   }
 
-  if (count * ncclTypeSize(datatype) > (size_t)rcclParamDdaLLTwoShotThreshold()) {
+  if (count * ncclTypeSize(datatype) > rcclDdaLLTwoShotThreshold(comm)) {
       return false;
   }
 
@@ -471,7 +470,7 @@ bool ddaLL128ArTwoShotEligible(ncclComm* comm, const void* sendbuff, void* recvb
     return false;
   }
 
-  if (count * ncclTypeSize(datatype) > (size_t)rcclParamDdaLL128TwoShotThreshold()) {
+  if (count * ncclTypeSize(datatype) > rcclDdaLL128TwoShotThreshold(comm)) {
     return false;
   }
 

@@ -13,6 +13,7 @@
 #include "debug.h"
 #include "algorithms/dda/dda_init_detail.h"
 #include "algorithms/dda/fabric/fabric_gpu_barrier.h"
+#include "rccl_common.h"
 
 #include <cuda_runtime.h>
 #include <hip/hip_ext.h>
@@ -40,8 +41,8 @@ static ncclResult_t ncclReduceScatterDdaFabricTyped(const void* sendbuff, void*
     return ncclInvalidArgument;
   }
 
-  // Use the block cap chosen at init (barrier flag buffer is sized for it).
-  const int nBlocksMax = comm->ddaFabricMaxBlocks;
+  // At most the block cap chosen at init (barrier flag buffer is sized for it).
+  const int nBlocksMax = rcclDdaFabricVmmBlocks(comm);
   // For reduce-scatter we use recvcount for grid calculation since each rank
   // processes its own shard.
   auto gridBlock = dda::common::getGridAndBlockDims(recvcount, sizeof(T), nBlocksMax);

@@ -1150,17 +1150,38 @@ ncclResult_t ncclCeAlltoAllv(struct ncclComm* comm, struct ncclCeCollArgs* args,
   uint8_t* myRecvBuff = (uint8_t*)args->recvBuff;
   void* peerRecvBuff;
   size_t offset, winOff;
-  size_t totalBytes = 0;
+  const int freq = (int)comm->ceColl.intraBatchSyncFreq;
+  bool roundSync = false;
   struct ncclCeBatchOpsParams batchOpsParams = {};
   NCCLCHECKGOTO(ncclCeInitBatchOpsParams(&batchOpsParams, comm->nRanks), ret, fail);
 
   NCCLCHECKGOTO(ncclAlltoAllvValidateSizeMatrix(args->sizes, comm->nRanks), ret, fail);
 
+  // ncclMemOpSync is collective: decide intra-batch sync from the full size matrix, not this rank's nonzero peers.
+  for (int src = 0; freq > 0 && src < comm->nRanks && !roundSync; src++) {
+    size_t* srcSizes = ncclAlltoAllvSendSizes(args->sizes, src, comm->nRanks);
+    int ops = 0;
+    size_t bytes = 0;
+    for (int dst = 0; dst < comm->nRanks; dst++) {
+      if (srcSizes[dst] != 0) {
+        ops++;
+        bytes += srcSizes[dst];
+      }
+    }
+    roundSync = ops > freq && bytes >= comm->ceColl.intraBatchSyncMsgThreshold;
+  }
+
   // Ensure all ranks are ready before starting transfers
   NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail);
 
   // Copy data to other ranks: send variable-sized chunk for each destination rank
   for (int r = 0; r < comm->nRanks; r++) {
+    // Sync per round of freq destination slots, so every rank issues the same number of barriers.
+    if (roundSync && r > 0 && r % freq == 0) {
+      NCCLCHECKGOTO(ncclCeLaunchBatchOps(comm, &batchOpsParams, stream, args), ret, fail);
+      batchOpsParams.numOps = 0;
+      NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail);
+    }
     int dstRank = (comm->rank + r) % comm->nRanks;
     const size_t chunkBytes = sendSizes[dstRank];
     if (chunkBytes == 0) {
@@ -1169,7 +1190,6 @@ ncclResult_t ncclCeAlltoAllv(struct ncclComm* comm, struct ncclCeCollArgs* args,
 
     uint8_t* srcPtr = mySendBuff + sendDispls[dstRank];
     uint8_t* dstPtr = myRecvBuff + recvDispls[comm->rank];
-    totalBytes += chunkBytes;
 
     if (dstRank == comm->rank) {
       // Local copy for own data
@@ -1197,13 +1217,14 @@ ncclResult_t ncclCeAlltoAllv(struct ncclComm* comm, struct ncclCeCollArgs* args,
     }
   }
 
-  // Check if we need to perform intra-batch synchronization
-  batchOpsParams.intraBatchSync =
-    (batchOpsParams.numOps > comm->ceColl.intraBatchSyncFreq && totalBytes >= comm->ceColl.intraBatchSyncMsgThreshold);
-
-  // Launch the batch operations
+  // Launch the last round; intraBatchSync stays false so ncclCeLaunchBatchOps adds no rank-local barriers
   NCCLCHECKGOTO(ncclCeLaunchBatchOps(comm, &batchOpsParams, stream, args), ret, fail);
 
+  // Graph replay needs an even barrier count so useCompletePtr returns to its pre-launch state.
+  if (roundSync && ncclCudaGraphValid(comm->planner.capturingGraph) && ((comm->nRanks - 1) / freq) % 2 != 0) {
+    NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail);
+  }
+
   // Ensure all transfers are complete across all ranks
   NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail);
 

@@ -3070,6 +3070,7 @@ static ncclResult_t topoGetAlgoInfo(struct ncclComm* comm, struct ncclTaskColl*
   }
   // Honor Tuner config if available
   if (!isTunerMatchFound && info->algorithm != NCCL_ALGO_PAT) {
+    rcclUpdateCollectiveAlgorithm(comm, nBytes, info, table);
     rcclUpdateCollectiveProtocol(comm, nBytes, info);
   }
   rcclSetPipelining(comm, nBytes, info);
@@ -3122,6 +3123,7 @@ static ncclResult_t topoGetAlgoInfo(struct ncclComm* comm, struct ncclTaskColl*
     }
     INFO(NCCL_TUNING, "post-adjustment based on threadThreshold:%i nBytes:%lu nc:%i", threadThreshold, nBytes, nc);
     rcclOverrideChannels(comm, info->func, nBytes, nc);
+    rcclUpdateCollectiveChannels(comm, nBytes, info, nc);
   }
 
 #ifdef ENABLE_ROCSHMEM
@@ -4765,8 +4767,8 @@ static ncclResult_t taskAppend(struct ncclComm* comm, struct ncclInfo* info) {
                                    opDev));
 
       } else if (!rcclNcclAlgoEnvIsSet() && !allGatherDecided && !alltoAllDecided &&
-                 rcclParamForceCe() && CeScratchAvailable && !hasSysmemSegment && comm->ddaScratch != nullptr &&
-                 recvBytes <= comm->ddaScratchBytes && info->coll != ncclFuncAllReduce) {
+                 rcclForceCeScratchEnabled(comm) && CeScratchAvailable && !hasSysmemSegment &&
+                 comm->ddaScratch != nullptr && recvBytes <= comm->ddaScratchBytes && info->coll != ncclFuncAllReduce) {
         INFO(NCCL_TUNING, "Using DDA scratch for CE collective, count=%zu, recvBytes=%zu", info->count, recvBytes);
         NCCLCHECK(ceCollTaskAppend(comm, info, /*sendWin=*/nullptr, /*recvWin=*/nullptr, comm->ddaScratch,
                                    comm->ddaPeerPtrsHost, opDev));

@@ -645,6 +645,9 @@ struct ncclPeerInfo {
   bool gpuCftMulticastSupport;
   bool gpuCftCountedSupport;
   uint32_t gitVersionHash;
+  // [RCCL] The rcclTuningPartition this rank asks for (RCCL_TUNING_PARTITION_MODE). The comm
+  // uses DPX only when every rank asks for it. Append later fields after this one.
+  int tuningPartition;
 };
 
 typedef enum ncclGroupTaskType {

@@ -79,8 +79,64 @@ struct rcclArchThresholds {
   // 0 means no suppression (symk may win at any size for that collective).
   size_t symMinR2[RCCL_DDA_FUNC_COUNT];
 
+  // Rank-scaled tier caps, bytes per rank. On a comm of at least
+  // rankScaledMinRanks ranks, a non-zero entry replaces the matching
+  // ddaLLMax / ddaLL128Max / ddaVmmMax total with nRanks * entry, clipped to
+  // rankScaledMaxBytes (which must not exceed the largest total above: fabric
+  // scratch is sized from the totals). Smaller comms, and entries of 0, keep the
+  // totals. The RCCL_DDA_*_THRESHOLD env vars still win over both.
+  size_t ddaLLMaxPerRank[RCCL_DDA_FUNC_COUNT];
+  size_t ddaLL128MaxPerRank[RCCL_DDA_FUNC_COUNT];
+  size_t ddaVmmMaxPerRank[RCCL_DDA_FUNC_COUNT];
+  int rankScaledMinRanks;     // 0 = no rank scaling
+  size_t rankScaledMaxBytes;
+
+  // AllReduce only: on comms of at least ddaArVmmFirstMinRanks ranks, DDA VMM
+  // takes unregistered messages in (ddaArVmmFirstMinBytes, nRanks *
+  // ddaArVmmFirstMaxPerRank] ahead of the LL tier, provided the VMM path is
+  // eligible. 0 ranks = off.
+  int ddaArVmmFirstMinRanks;
+  size_t ddaArVmmFirstMinBytes;
+  size_t ddaArVmmFirstMaxPerRank;
+
+  // DDA VMM (AllReduce / ReduceScatter) launches min(init block cap,
+  // clamp(ddaVmmBlockBudget / nRanks, ddaVmmMinBlocks, ddaVmmMaxBlocks)) blocks
+  // on comms of at least rankScaledMinRanks ranks. Budget 0 = the init block cap.
+  int ddaVmmBlockBudget;
+  int ddaVmmMinBlocks;
+  int ddaVmmMaxBlocks;
+
+  // Ring / Tree crossovers that replace the tuning model's ranges on single-node
+  // comms (nNodes == 1). Sizes are per rank as rcclGetSizePerRank() reports them
+  // (AllReduce: the whole message, hence nRanks x the entry). 0 = keep the
+  // model's range. The two AllReduce LL128 entries apply at any rank count, the
+  // rest from rankScaledMinRanks ranks.
+  size_t arRingLL128MaxPerRank;   // Ring AllReduce LL128 up to nRanks x this
+  bool arTreeLL128Unlimited;      // Tree AllReduce keeps LL128 at every size
+  size_t arTreeMaxPerRank;        // AllReduce runs Tree instead of Ring up to nRanks x this,
+  size_t arTreeLLMaxPerRank;      // with LL up to nRanks x this and LL128 above
+  size_t agRingLL128MaxPerRank;   // AllGather Ring LL128 up to this per rank
+  size_t rsRingLL128MinPerRank;   // ReduceScatter Ring LL128 above this per rank,
+  size_t rsRingLL128MaxPerRank;   // up to this (2- and 4-byte types)
+  size_t rsRingLL128MaxPerRank1B; // or this (1-byte types)
+  size_t rsRingSimpleCappedMaxPerRank;  // Ring Simple ReduceScatter above the LL128 band,
+  int rsRingSimpleCappedChannels;       // up to this per rank, runs on at most this many channels
+
+  // Force CE-Scratch for unregistered AllGather / AlltoAll / Gather / Scatter
+  // whenever the receive fits ddaScratch, on top of the ceNonRegMin/Max window.
+  // NCCL_CTA_POLICY_ZERO forces it on every arch (it is the one zero-CU path for
+  // unregistered buffers). RCCL_FORCE_CE=0 turns it off everywhere. A null table forces it.
+  bool forceCeScratch;
+};
+
+// Compute-partition mode a gfx1250 comm is tuned for (RCCL_TUNING_PARTITION_MODE). It picks
+// the table only; other archs have one table for every mode.
+enum rcclTuningPartition {
+  RCCL_TUNING_PARTITION_SPX = 0,
+  RCCL_TUNING_PARTITION_DPX = 1,
 };
 
-const rcclArchThresholds* rcclGetArchThresholds(const char* gcn);
+const rcclArchThresholds* rcclGetArchThresholds(const char* gcn,
+                                                rcclTuningPartition partition = RCCL_TUNING_PARTITION_SPX);
 
 #endif  // RCCL_ARCH_THRESHOLDS_H_