VadimKutovoi · Nov 3, 2026 · Oct 16, 2026 · Oct 16, 2026 · Oct 16, 2026 · Oct 16, 2026
@@ -48,6 +48,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * 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_TUNER_EMBEDDED_CONFIG` (default `1`): the built-in CSV tuner configurations in `tuner/*.csv` are now compiled into `librccl.so` and used as the stock tuning defaults, so a cluster no longer has to stage the CSV on every node or on NFS. Set `0` to ignore the embedded configurations. `NCCL_TUNER_CONFIG_FILE` and the on-disk tuner directories still take precedence, and `NCCL_TUNER_PLUGIN=none` still disables the tuner entirely.
 * `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.
+* `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`.
 * `RCCL_CE_AR_2SHOT_MAX_BYTES` (default `-1`): overrides the 2-shot AllReduce size cap from the arch table. Set to a positive value to override `ceNonRegMax[AR]`. Renamed from `RCCL_CE_AR_MAX_MSG_BYTES`.
 * `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.
 * `RCCL_CE_AR_REG_MAX_MSG_BYTES` (default `-1`): overrides the registered CE AllReduce size cap from the arch table.

@@ -667,6 +667,24 @@ in the following table.
       - | ``0``: Disabled (default).
         | ``1``: Enabled.
 
+    * - | ``RCCL_IB_UDMA_POLICY``
+        | Selects how queue pairs (QPs) are assigned to the two uDMA pipelines
+          of an AINIC. Applies only to AINIC with the ``IB-CAST`` transport.
+          With QP sharing enabled, the sharing groups choose the pipeline in
+          place of the per-channel assignment of ``0`` and ``1``.
+      - | ``0``: Pin each channel's send QPs, and separately its receive QPs,
+          to one pipeline per device, alternating pipelines as each is first
+          used (default).
+        | ``1``: Round-robin P2P QPs across the pipelines in creation order.
+          Collective QPs behave as with ``0``.
+        | ``2``: Leave the pipeline choice to the driver.
+        | ``3``: Split each connection across both pipelines. The default
+          ``NCCL_IB_QPS_PER_CONNECTION`` of 2 is enough. Splitting is lost with
+          ``NCCL_IB_QPS_PER_CONNECTION=1`` or ``RCCL_IB_QPS_PER_P2P=1``, and on
+          RMA comms, which use one QP; every QP of a channel then uses the same
+          pipeline.
+        | Other values log a warning and fall back to ``0``.
+
     * - | ``NCCL_NCHANNELS_PER_NET_PEER``
         | Sets the number of channels used per network (remote) peer.
         | This overrides the value of the ``nChannelsPerNetPeer`` field in

@@ -74,18 +74,34 @@ enum ncclIbRequestMatchingScheme {
 // sharing enabled, the sharing-group alternation replaces the channel affinity
 // of ChannelAffinity, and of collective QPs under P2pRoundRobin.
 enum struct ncclIbCastUDMAPolicy : int {
-  // Pin every QP of a channel to one pipeline, chosen alternately as channels are first seen
+  // Pin each channel's send QPs, and separately its receive QPs, to one pipeline
+  // per device, chosen alternately as each is first seen
   ChannelAffinity = 0,
   // Round-robin P2P QPs in creation order; collective QPs keep channel affinity
   P2pRoundRobin = 1,
   // Leave the PD mask unchanged so the driver picks the pipeline
   DriverDefault = 2,
   // Split each connection across both pipelines: channel-id popcount parity picks the
-  // starting pipeline and consecutive per-device QP indexes alternate
+  // starting pipeline and consecutive per-device QP indexes alternate. The default
+  // NCCL_IB_QPS_PER_CONNECTION of 2 is enough. With one QP per device per connection
+  // (NCCL_IB_QPS_PER_CONNECTION=1, or RCCL_IB_QPS_PER_P2P=1 for P2P) every QP has
+  // qpIdx 0 and this reduces to channel-id popcount parity. RMA comms also use one QP.
   SplitConnection = 3,
+  // New values go before Count; First/Last bound the RCCL_IB_UDMA_POLICY range
+  // check, so update the static_asserts below when adding one.
+  Count,
   First = ChannelAffinity,
-  Last = SplitConnection,
+  Last = Count - 1,
 };
+static_assert(ncclIbCastUDMAPolicy::First == ncclIbCastUDMAPolicy::ChannelAffinity &&
+                (int)ncclIbCastUDMAPolicy::First == 0,
+              "ncclIbCastUDMAPolicy must start at ChannelAffinity = 0");
+static_assert(ncclIbCastUDMAPolicy::Last == ncclIbCastUDMAPolicy::SplitConnection,
+              "ncclIbCastUDMAPolicy::Last must be the last policy; update this assert when adding one");
+
+static inline bool ncclIbCastUDMAPolicyIsValid(int64_t policy) {
+  return policy >= (int64_t)ncclIbCastUDMAPolicy::First && policy <= (int64_t)ncclIbCastUDMAPolicy::Last;
+}
 
 struct ncclIbMr {
   uintptr_t addr;

@@ -83,7 +83,7 @@ constexpr int8_t ncclIbCastUDMANone = (int8_t)-1;
 
 static ncclIbCastUDMAPolicy ncclIbCastGetUDMAPolicy() {
   int64_t policy = rcclParamIbCastUDMAPolicy();
-  if (policy < (int64_t)ncclIbCastUDMAPolicy::First || policy > (int64_t)ncclIbCastUDMAPolicy::Last) {
+  if (!ncclIbCastUDMAPolicyIsValid(policy)) {
     return ncclIbCastUDMAPolicy::ChannelAffinity;
   }
   return (ncclIbCastUDMAPolicy)policy;
@@ -98,6 +98,8 @@ static int8_t ncclIbCastSelectUDMA(const struct ncclIbQpCreateAttr* createQpAttr
   uint8_t udmaCount = IbCastAinicRoce ? IbCastDevs[createQpAttrs->ibDevN].capsProvider.ionic.udmaCount : 2;
 
   if (policy == ncclIbCastUDMAPolicy::SplitConnection) {
+    // qpIdx only advances with >= 2 QPs per device per connection; otherwise
+    // it is always 0 and only the channel-id parity varies.
     uint8_t udmaEngine = (__builtin_parity(static_cast<unsigned>(createQpAttrs->channelId))
                          + createQpAttrs->qpIdx) % udmaCount;
     return (int8_t)(1u << udmaEngine);

@@ -685,7 +685,7 @@ ncclResult_t IbCastInitDevices(ncclDebugLogger_t logFunction, ncclProfilerCallba
            IbCastUseInline ? "Enabled" : "Disabled", IbCastGdrFlushDisable ? "Disabled" : "Enabled");
 
       int64_t udmaPolicy = rcclParamIbCastUDMAPolicy();
-      if (udmaPolicy < (int64_t)ncclIbCastUDMAPolicy::First || udmaPolicy > (int64_t)ncclIbCastUDMAPolicy::Last) {
+      if (!ncclIbCastUDMAPolicyIsValid(udmaPolicy)) {
         WARN("NET/IB : RCCL_IB_UDMA_POLICY=%ld is invalid (must be %d-%d). Using RCCL_IB_UDMA_POLICY=%d.", udmaPolicy,
              (int)ncclIbCastUDMAPolicy::First, (int)ncclIbCastUDMAPolicy::Last,
              (int)ncclIbCastUDMAPolicy::ChannelAffinity);

@@ -30,6 +30,12 @@ ncclResult_t IbCastPortRecoveryQpsRestoreAinic(struct ncclIbPortRecoveryContext*
     }
     int devIndex, qpIdxInDev;
     NCCLCHECK(IbCastQpStripe(qpIndex, recoveryContext->resCtx->baseComm->vProps.ndevs, &devIndex, &qpIdxInDev));
+    if (devIndex != localQp->devIndex) {
+      WARN("NET/IB: %s: QP %d has devIndex %d but striping over %d devices places it on device %d (comm=%p)", __func__,
+           qpIndex, localQp->devIndex, recoveryContext->resCtx->baseComm->vProps.ndevs, devIndex,
+           recoveryContext->resCtx->baseComm);
+      return ncclInternalError;
+    }
 
     uint32_t oldQpn = localQp->qp->qp_num;
     INFO(NCCL_NET, "NET/IB: %s: Destroying QP %d on device %d (comm=%p, old_qp_num=%u)", __func__, qpIndex,