Repository navigation
fix(rccl): RCCL_IB_UDMA_POLICY review follow-ups - #13727
Merged
VadimKutovoi merged 6 commits intoNov 3, 2026
Merged
Conversation
github-actions
Bot
added
documentation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
VadimKutovoi
marked this pull request as ready for review
October 20, 2026 15:51
VadimKutovoi
requested review from
a team,
alex-breslow-amd and
amd-arozanov
and
a balanced review from Copilot
October 20, 2026 15:51
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
Hardware validation of the changed one-QP balancing behavior remains pending on MI355X/AINIC.
No open issues 🎉
What changed in this PR
Follows up on #12948 to refine RCCL’s AINIC uDMA policy handling while preserving default behavior.
Changes:
- Separates send and receive round-robin counters to balance odd P2P QP counts.
- Adds enum guards, shared policy validation, and a recovery device-index check.
- Documents policy options and the minimum QP count for connection splitting.
| File | Description |
|---|---|
| projects/rccl/src/transport/net_ib_cast/p2p_resiliency_recovery_ainic.cc | Checks device-index consistency during recovery. |
| projects/rccl/src/transport/net_ib_cast/init.cc | Uses shared policy validation. |
| projects/rccl/src/transport/net_ib_cast/connect.cc | Separates send/receive counters and reuses validation. |
| projects/rccl/src/transport/net_ib_cast/common_cast.h | Adds enum guards, validation helper, and policy comments. |
| projects/rccl/docs/api-reference/env-variables.rst | Documents uDMA policy configuration. |
| projects/rccl/CHANGELOG.md | Announces the policy variable. |
Review effort: Balanced
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
ARGUS
Findings
1 more finding(s), not posted inline
projects/rccl/docs/api-reference/env-variables.rst:606(Low) env-variables.rst: policy 0 says all QPs of a channel share one pipeline, but send and receive QPs are assigned separately
Bottom line: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
VadimKutovoi
force-pushed
the
users/vkutovoi/ROCM-31118-udma-policy-followup
branch
from
October 20, 2026 16:22
0316f1b to
b4a8808
Compare
amd-arozanov
requested changes
Oct 26, 2026
alexander-sannikov
approved these changes
Oct 30, 2026
alexander-sannikov
left a comment
alexander-sannikov
left a comment
Contributor
There was a problem hiding this comment.
LGTM, in case doc comments from @amd-arozanov are addressed.
Derive Last from a trailing Count enumerator and static_assert that First == ChannelAffinity == 0 and Last == SplitConnection, so adding a policy without updating the range check fails to compile. No enum value is used in a switch, so Count does not trigger -Wswitch.
Add ncclIbCastUDMAPolicyIsValid() to common_cast.h and use it for both the init-time warning and ncclIbCastGetUDMAPolicy(), so the two checks cannot drift apart. No behavior change.
IbCastPortRecoveryQpsRestoreAinic recomputes qpIdx from the QP stripe but discarded the stripe's devIndex. Fail with ncclInternalError if it disagrees with the QP's recorded devIndex, since the restored qpIdx would then be wrong. Connect-time QPs are always striped this way and resiliency is off under QP sharing, so this does not fire in practice.
With one QP per device per connection every QP has qpIdx 0, so RCCL_IB_UDMA_POLICY=3 only varies by channel-id popcount parity and behaves like a channel affinity. Say so in the enum and at the selection site, and name the knobs that raise the QP count. Comment only.
Add a CHANGELOG entry and an env-variables page entry for the AINIC uDMA pipeline policy: IB-CAST only, values 0-3, default 0, out-of-range values warn and fall back to 0, and policy 3 needs at least two QPs per device per connection.
VadimKutovoi
force-pushed
the
users/vkutovoi/ROCM-31118-udma-policy-followup
branch
from
November 2, 2026 15:41
b4a8808 to
800ffe9
Compare
Policy 0 pins send and receive QPs separately per device, policy 1 round-robins P2P QPs on one counter, and the default of two QPs per connection already satisfies the split. Co-authored-by: Cursor <cursoragent@cursor.com>
amd-arozanov
approved these changes
Nov 2, 2026
Contributor
pr-rocm-systems-13727-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-13727-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
alex-breslow-amd
approved these changes
Nov 2, 2026
alex-breslow-amd
left a comment
alex-breslow-amd
left a comment
Contributor
There was a problem hiding this comment.
Thank you for completing this follow-up. :-)
VadimKutovoi
deleted the
users/vkutovoi/ROCM-31118-udma-policy-followup
branch
November 3, 2026 08:51
pmoutsias-amd
added a commit
that referenced
this pull request
Nov 3, 2026
…the RCCL 2.32.3 for ROCm 10.3.0 (Unreleased) section (#14929) This PR does two things to `projects/rccl/CHANGELOG.md`. 1. Removes the duplicate `## RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` heading. #13544 removed it, and #13128 brought it back. The section is not a pure copy: 48 of its entries have no identical copy elsewhere. 47 of them are folded into `## RCCL 2.32.3 for ROCm 10.2.0` instead of being deleted, and the 48th is the older wording of the `RCCL_CE_AR_STAGING_BYTES` entry, which is dropped in favor of the newer one. Where the same entry appeared in both sections, one copy is kept. 2. Moves the entries below, which sit in release sections but whose code was merged to develop after the `rocm-systems` commit that TheRock `release/therock-10.2` pins (`ded49fd`) and is not in that commit. This moves 53 of them into `## RCCL 2.32.3 for ROCm 10.3.0 (Unreleased)`, under the matching subsections. Not moved: 10 entries that were attributed to #13544 (the CHANGELOG-only PR that last touched them), so the code PR could not be traced. They stay in the 10.2.0 section and are unverified, so the RCCL owners should confirm them. They begin: `RCCL_TDM_LL128_ENABLE`, `Ll128TdmStoreTest`, the TDM reduction path for SIMPLE, Anvil SDMA strong GIN signals, VA GIN signals on Anvil SDMA, `NCCL_P2P_DISABLE=1` on shared topology nodes, queue evictions at process exit, multi-receive `irecv` with `NCCL_IB_USE_INLINE=1`, Anvil SDMA signal cells across logical contexts, and GIN signal shadow sizing. - #13328 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_BCM_LINKS_PATH`: fallback directory path for Broadcom PEX890xx multi-host switch inter-switch-link topology. On systems where a single physical PEX switch is partitioned across multiple CPU domains (upstream ports), RCCL needs to know which ports share the same switch so it can classify cross-partition PCIe paths as `PXB` instead of `PHB` and enable GDR across them. The `switch_discovery` kernel module exposes this via sysfs (`/sys/kernel/pci_switch_link/virtual_switch_links`); when that path is not available, RCCL reads the same directory layout from `RCCL_BCM_LINKS_PATH`. The companion userspace tool `rccl_bcm_isl_discover.sh` discovers the switch groups by reading the PCIe Device Serial Number (DSN) from extended config space (vendor `1000`, device `c030`) and populates this directory. - #12742 (merged 2026-10-18, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_MLOPART`: allows GIN on MLO/CPX partitions. A partition reaches xGMI and the network over the physical device's paths, so GIN is as available to a partition as to the whole GPU. By default a single-node partitioned communicator keeps GIN, and one spanning several nodes does not, because past one node admitting a partition also reaches the hierarchical Copy Engine path. Set `1` to allow GIN on partitions on any number of nodes, or `0` to restore the previous partition exclusion. - #3894 (merged 2026-10-16, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_TREE_COMPACT` (default `0`): builds the inter-node trees with a half-interleave node remapping, for jobs whose nodes are ordered by network domain. Both trees keep cross-domain hops at the leaves only when the nodes split evenly into two domains; with an uneven split, the first tree does so when the first domain holds at least half of the nodes, and the second tree still crosses domains. All ranks in a communicator must use the same value; initialization fails with `ncclInvalidUsage` otherwise. - #12856 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_FABRIC_TELEMETRY_ENABLE` (default `0`): samples every amd-smi fabric telemetry category for the GPUs this process opens communicators on and logs the counters that moved, plus a closing report over the whole sampled span when a device's last communicator is destroyed. Diagnostic only, and off by default. Also needs `RCCL_USE_AMD_SMI_LIB=1` and `render`-group access to the IFoE config character device; without either, RCCL logs why telemetry was skipped and initialization continues. - #12856 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_FABRIC_TELEMETRY_INTERVAL_MS` (default `1000`): sampling period for `RCCL_FABRIC_TELEMETRY_ENABLE`, clamped to a 100 ms floor and a one-hour ceiling. A non-positive value disables telemetry entirely. - #12573 (merged 2026-10-14, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Hierarchical symmetric `ncclAllGather` for multi-node using GIN and LSA (`AllGather_RailRing_LsaST`): a GIN ring over the rails team moves data between nodes and LSA stores broadcast it to the other GPUs in each node. - #14140 (merged 2026-11-01, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Hierarchical symmetric `ncclAllGather` for partition mode in gfx1250 scale-up domain using LSA (`AllGater_HierLsaST`): each rank stores its data once to its counterpart partition on every other GPU, which forwards it to the other partitions there, so data crosses to each GPU once instead of once per partition. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_ALLREDUCE_GFX1250_ENABLE` (default `0`): opt in to GinScatter AllReduce on gfx1250 for symmetric sum messages from 64 MiB through 4 GiB. Left off, gfx1250 keeps the pre-existing AllReduce selectors. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_GIN_ALLREDUCE_SCRATCH_BYTES` (default 512 MiB): GinScatter incoming scratch. `0` skips the allocation and GinScatter is not selected. The buffer is allocated only on the GinScatter path. - #13416 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `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. - #13727 (merged 2026-11-03, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `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`. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `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. - #12310 (merged 2026-10-28, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_REDUCESCATTER` (default `0`): opt-in Copy Engine ReduceScatter on a single-node symmetric communicator with `CTA_POLICY_ZERO`. Always stages through `ceARTmpBuf`; user-buffer registration is not required. No message-size cap (large messages are pipelined). - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_COLL_MAX_BYTES` (default `-1`): overrides the CE registered-window and 2-shot size cap for all non-AllReduce collectives (AllGather, AlltoAll, etc.). When `-1`, resolved from `ceRegMax[func]` or `ceNonRegMax[func]` in the per-arch table. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_CE_COLL_MIN_BYTES` (default `-1`): overrides the minimum message size for CE non-AllReduce paths (both registered-window and scratch). When `-1`, resolved from `ceNonRegMin[func]` in the per-arch table. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_SYM_K_MAX_BYTES` (default `-1`): overrides the symmetric-kernel upper bound for all collectives. When `-1`, resolved from `symMaxR2[func]` (eager) or `symMaxR2Graph[func]` (graph-capture) in the per-arch table. `kThreshUnlimited` (SIZE_MAX) in the table means no suppression. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_SYM_K_MIN_BYTES` (default `-1`): overrides the symmetric-kernel lower bound for all collectives. When `-1`, resolved from `symMinR2[func]` in the per-arch table. Below this size, DDA is preferred over symk. - #13995 (merged 2026-10-27, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_GDR_FLUSH_FLAG` (default `0`): replaces the GPUDirect RDMA flush of Simple-protocol receives in the `IB` network transport. Instead of the proxy fencing each receive with an RDMA read and waiting for it, the sender writes a per-QP flag into GPU memory after the payload without relaxed ordering, and the GPU waits for it. A communicator uses it only when every rank enables it and can use it (gfx942 or gfx950, the `IB` transport, and no adaptive routing, out-of-order receive queues, ID-based receive matching (`NCCL_IB_RECEIVER_SIDE_MATCHING_SCHEME=1`) or port-failover resiliency (`NCCL_IB_RESILIENCY_PORT_FAILOVER`)); otherwise every rank keeps the RDMA read flush. Connections over a fused NIC (NIC Fusion, including a fused vNIC chosen by subnet-aware routing), with more than 8 QPs per request, or that cannot register the flag memory also keep it. On gfx950 with 2 to 8 nodes of 8 GPUs, such a communicator also selects Ring/LL128 from 16 MB per node, Tree/Simple from 32 MB per node and Ring/Simple from 64 MB per node for `ncclAllReduce`, and switches `ncclReduceScatter` from LL128 to Simple above 6 MB per rank instead of 8 MB, unless `NCCL_ALGO` or `NCCL_PROTO` is set. On MI355X it raises large `ncclAllReduce` and `ncclReduceScatter` bandwidth by 10 to 22% and 1 to 8 MB `ncclAlltoAll` by 4 to 12%. - #13995 (merged 2026-10-27, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_IB_GDR_FLUSH_FLAG_TUNE` (default `0`): additional tuning for communicators that use `RCCL_IB_GDR_FLUSH_FLAG` on gfx950 with 8 GPUs per node on 2 to 8 or 16 nodes. Ring and tree connections get a 4x LL128 buffer (send/recv keeps the default, and `NCCL_LL128_BUFFSIZE` disables it); `ncclAllReduce` selects Ring/LL128 from 12 MB per node up to what one ring loop carries, then Tree/Simple up to 64 MB per node and Ring/Simple on 48 channels up to 128 MB per node, and Tree/Simple from 8 to 64 MB per node on 16 nodes; `ncclReduceScatter` keeps LL128 up to 7 MB per rank, then Simple, on 16 nodes too. All ranks must use the same value. - #13757 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_PCIE_ORDERED_RING` (default `1`, also accepted as `NCCL_PCIE_ORDERED_RING`): pins the ring graph search to the order GPUs appear in the topology array, so every channel uses the same ring. Without it the search picks a different permutation per channel, so a device exchanges with many different peers; pinning the order reduces fan-out to one peer per direction per device across all rings, which cuts the connections, FIFO buffers and setup work each device carries. On gfx1250 partitioned (DPX) GPUs the unpinned search also mixes scale-up and scale-in hops between channels; pinning keeps the two partitions of a die adjacent, though ring throughput is bounded by its slowest hop either way. This constrains GPU visitation order only: path validation, bandwidth constraints and backtracking are unchanged. It applies to gfx1250 ring graphs on a single host where every rank is a GPU of that topology, so trees, CollNet, other architectures and multi-host communicators are unaffected, and `NCCL_GRAPH_FILE` and `NCCL_RINGS` still take precedence. Set `0` to restore the unconstrained search, or `2` to fail communicator initialization instead of falling back when no ordered ring passes validation. - #10885 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `NCCL_MULTI_SEGMENT_REGISTER` default changed `1` -> `0` on HIP and stays `1` on other builds. On HIP, multi-segment registration stays off until a job sets the variable. Set `1` to register one DMA-BUF memory region per physical segment where the transport supports it. - #10885 (merged 2026-10-26, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * NET registration counts physical segments over the whole registered buffer, not only the range a collective touches. With `NCCL_MULTI_SEGMENT_REGISTER=0`, a collective on any part of a multi-segment registration skips NET registration and stages, instead of retrying a single-region registration on every collective. - #10891 (merged 2026-10-25, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * GIN window registration on HIP fails with `ncclInvalidUsage` when the window spans several device-memory segments, instead of registering a memory region that covers only the first segment, and when it contains host VMM memory and `NCCL_DMABUF_ENABLE=0`. - #12947 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `install.sh --all_unrolls` (`-DBUILD_ALL_UNROLLS=ON`) now adds the unroll factors the multi-arch default skips (8 and 16) on top of each architecture's existing set. It no longer replaces that set or clears the gfx1250 pin on unroll 32. The flag stays off by default, so a normal build does not compile unroll 8 or 16. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The GIO topology model is applied only to systems with exactly 16 GPUs. It was previously applied to any system with 16 or more GPUs, where the 16-GPU model does not describe the topology; such systems now fall through to the other models. - #13581 (merged 2026-10-19, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Enabled the single-node one-slice optimization for gfx1250, which was previously limited to gfx942 and gfx950. Ring Simple AllGather, AllReduce and ReduceScatter now use one slice per chunk instead of two, halving the per-chunk peer handshakes. The gain is largest below about 128 MB and tapers to a few percent at multi-GB sizes, where the ring is already near the link bandwidth. LL, LL128 and non-ring algorithms are unaffected. - #14123 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * On gfx1250, the DDA LL128 `ncclAllGather`, `ncclAllReduce` and `ncclReduceScatter` receivers keep each line's payload from the cooperative load that saw its flag instead of reading the slice a second time. The DDA scratch bank is now floored to 128 bytes instead of 16, so an LL128 wire line never straddles two 128-byte lines when `RCCL_DDA_FABRIC_BUFFER_SIZE` sets an unaligned size. Results are unchanged. Other architectures are unaffected. - #13415 (merged 2026-10-19, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `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. - #13406 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * 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. - #13357 (merged 2026-10-23, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `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. - #13357 (merged 2026-10-23, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * 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. - #12291 (merged 2026-10-21, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * The AMD Pensando 'ionic' RDMA provider for the rocSHMEM GDA GIN backend (`NCCL_GIN_TYPE=6`) now properly enforces the required support for creation of 'collapsed' single-entry completion queues. If network driver and firmware support for this feature is unavailable, an initialization failure will occur rather than a silent fallback to plain completion queues. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * `RCCL_DDA_LL128_RS_THREADS` removed (was default `1024`): reduce-scatter LL128 now launches a fixed 512-thread block, one warp per slice, so the thread count is no longer a tuning knob. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, causing an `ncclGather` performance regression on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed LL-protocol latency on gfx950 when RCCL is built against a ROCm older than 7.0.2 without the 128-bit global load/store builtins (for example ROCm 7.0.1, or `-DDWORDX4_INTRINSICS=OFF`). Every store to the user output buffer was followed by a release fence and an L2 writeback, which made 4 KB to 64 KB `ncclBroadcast` about 30% slower on MI355X. Builds against ROCm 7.0.2 or later were not affected and are unchanged. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed a 15 to 30% `ncclScatter` bandwidth loss from 16 MB upward on multi-node AINIC jobs, caused by the move to two network channels per peer (for example on 4 nodes of 8 GPUs). Scatter now uses and connects the most channels per peer whose flows still fit the P2P channel pool. Point-to-point operations and the other collectives keep their channel count. - #13128 (merged 2026-11-02, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.32.3 for ROCm 10.2.0` section: > * Fixed `ncclCommAbort` and `ncclCommRevoke` hanging when a DDA collective (IPC on gfx942/gfx950, fabric LL/LL128/VMM on gfx1250) was waiting on a peer that never joined. DDA kernels now observe the communicator abort flag while waiting. A parent communicator shrunk with `NCCL_SHRINK_ABORT` no longer uses DDA for later collectives. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * DDA LL128 eligibility functions (`ddaLL128ArOneShotEligible`, `ddaLL128ArTwoShotEligible`) no longer check `RCCL_DDA_LL`. LL and LL128 are now independently gated: `RCCL_DDA_LL=0` disables only the LL tier; `RCCL_DDA_LL128=0` disables only the LL128 tier. Both checks are enforced via `rcclDdaLLThresholdTab` / `rcclDdaLL128ThresholdTab` at the selector level, which return 0 when the respective param is disabled. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * DDA LL128 is disabled for ReduceScatter and AlltoAll on gfx1250: `ddaLL128Max[RS]` and `ddaLL128Max[A2A]` are 0 in the gfx1250 arch table. - #13005 (merged 2026-10-15, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * With `NCCL_CTA_POLICY=2` (`NCCL_CTA_POLICY_ZERO`), AllReduce on registered symmetric windows now takes the registered CE path whenever it is available, even when the symmetric kernel is eligible, matching AllGather and upstream NCCL. A non-zero arch-table `ceRegMax[AllReduce]` no longer caps it under this policy; a zero entry and an explicit `RCCL_CE_AR_REG_MAX_MSG_BYTES` still apply. - #12386 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_CE` (default `1`): acts as a prerequisite gate for CE-scratch selection. When `RCCL_FORCE_CE=0`, CE-scratch is disabled entirely for all collectives regardless of other flags. When non-zero, CE-scratch fires if `CTA_POLICY_ZERO` (`-x 2`) is set (bypassing the size window) or if the message size falls within the arch-table `ceNonRegMax` window. AlltoAll CE-scratch now has a size window (`ceNonRegMax[AlltoAll]` in the per-arch table, currently 0 on all arches). - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Registered-window CE AllReduce is now eligible without requiring ``NCCL_CTA_POLICY=2`` > (``CTA_POLICY_ZERO``); it is gated only by the per-arch table window and ``RCCL_CE_ALLREDUCE``. > ``RCCL_FORCE_CE_ALLREDUCE`` now applies exclusively to the non-registered (2-shot/staged) CE > AllReduce path. - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_SYMMETRIC` (default `0`): when set to a non-zero value, biases the `ncclAllReduce` > backend selector to prefer the symmetric kernel path over all other candidates. The selector > still requires the operand buffers to be registered symmetric windows; if the symmetric > path is ineligible the selector falls back normally. - #13098 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * `RCCL_FORCE_DDA` (default `0`): when set to a non-zero value, biases the `ncclAllReduce` > backend selector to prefer DDA (Direct Data Access) paths over all other candidates, > including symmetric and Copy Engine. If no DDA path is eligible for the given message > size and communicator the selector falls back normally. - #14254 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed an intermittent hang, and possible wrong results, in the first DDA fabric LL all-reduces after communicator initialization on gfx1250 (`NCCL_MNNVL_ENABLE=1`). Initialization zeroed the LL scratch buffer and epoch counters with an asynchronous null-stream memset that could land after a rank's first LL kernel, resetting its epoch so that rank used a different flag and scratch bank than its peers. Initialization now waits for the zero fills, and the ranks agree that every rank zeroed its buffers before any rank can launch an LL kernel. If any rank fails to set up its LL buffers or GPU barrier, every rank disables the fabric DDA path together. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed memory corruption, uninitialized reads and a division by zero when parsing malformed or oversized `NCCL_RINGS` values. A ring list with more channels than `MAXCHANNELS`, a ring naming more GPUs or more NET endpoints than supported, or a GPU or NET number too large for an `int` wrote past the channel and parser arrays or wrapped around; NET slots a ring did not name were copied into the graph uninitialized; and input from which no channel could be parsed divided the system bandwidth by zero channels. Such input is now rejected with `ncclInvalidUsage`, and bandwidth is zero when no channel is parsed. Parsing a `NULL` ring string returns `ncclInvalidArgument` instead of dereferencing it. A ring that names its NETs between its GPUs rather than at its ends, for example `"0 N0 1 2 N1 3"`, no longer stores a negative NET id in the graph, which failed multi-node initialization; such a channel is ignored. - #13509 (merged 2026-10-31, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed a crash and silent misconfiguration when parsing malformed `RCCL_TREES` values. A tree not terminated by `)`, for example `"(1)(2)(3)(4)(5)(6"`, searched past the end of the string and segfaulted; a GPU number larger than an `int` wrapped around and built a tree over an unintended GPU, as `"(4294967299)"` did over GPU 3; and a stray non-digit character resolved to a GPU number left over from the previous tree, so `"(5)(6)|(0)(x)"` parsed as `"(0)(6)"`. The number of trees, the GPUs per tree, the transcribed tree length and the GPU index used for remapping are now bounded. Invalid input is rejected with `ncclInvalidUsage`, and a `NULL` tree string returns `ncclInvalidArgument`. - #13815 (merged 2026-10-30, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed wrong results and hangs in application kernels that use the GIN device API, such as the `07_kernel_fusion` GIN and hybrid RMSNorm examples. `nccl_device/coop.h` selected `__syncwarp` from `ROCM_VERSION`, which only RCCL's own build defines, so in application code the thread and warp cooperative groups synchronized the whole block with `__syncthreads()`. The GIN barrier calls them from a subset of threads, which left the block's waves out of step. The headers now fall back to the HIP version when `ROCM_VERSION` is not defined. - #12743 (merged 2026-10-20, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed the GIN plugin never being finalized for a communicator that had connected a GIN backend, leaking the plugin's per-communicator context and leaving its reference count one too high after every `ncclCommDestroy`. `ncclGinHostFinalize` closed the GIN connections and then zeroed the whole `ncclGinState`, so the `ncclGinFinalize` that follows it in `commFree` saw no active backend and skipped the plugin `finalize()` call. Host finalize now clears only the per-connection state and hands the backend records to `ncclGinFinalize`, which finalizes each one on the last shared-resource reference. `commFree` returns the first of a `closeColl` or plugin-finalize error after those frees; if both fail, the `closeColl` error is the one reported. An `ncclRmaProxyFinalize` error still aborts the rest of teardown immediately. On the Anvil SDMA backend (`NCCL_GIN_TYPE=7`) this was one leaked init context per connected communicator. - #12679 (merged 2026-10-16, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed a crash (`HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`) on gfx110x and gfx120x when the Tree algorithm runs over shared memory with more channels than the Tree graph search found, for example with `NCCL_ALGO=Tree` and `NCCL_P2P_DISABLE=1`. Channels added after `ncclTopoPreset` had no intra-node tree links; `ncclTopoPostset` now sets the tree and CollNet chain links on every channel, as it already did for rings. - #11133 (merged 2026-10-14, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed the host one-sided RMA API being unavailable on multi-node communicators whose ranks had fused NICs. NIC fusion merges devices for the regular transport only; GIN and RMA select from their own topology node types and never receive a fused device. - #12750 (merged 2026-10-18, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.31.2 for ROCm 10.2.0 (Unreleased)` section: > * Fixed Anvil SDMA GIN small-message AllToAll latency after `ncclGinFenceLevel` support. A sub-threshold `gin.put`, a `gin.putValue`, or a standalone `gin.signal` emitted two system-scope fences per peer: one in `fenceBeforeSignal` and one in `signalPeer`. `signalPeer` now emits a single system-scope release immediately before the remote atomic, rather than an agent-scope release before it and a system-scope fence after it, which also stops payload stores from passing the signal cell. A clean-queue IPC `SignalInc` skips `fenceBeforeSignal` entirely, since it would only repeat that release. The SDMA `quiet()` before a signal now runs only when the call itself issued SDMA; previously any call carrying a signal or counter quieted the peer queue, including an IPC put after an earlier SDMA put to the same peer, which now relies on `Flush` to drain that queue. The `NCCL_GIN_ANVIL_SDMA_IPC_AGENT_FENCE` debug knob is removed. Each `Flush` thread drains the peers it owns and then clears exactly the bits it drained, over a mask disjoint from every other thread's. Previously `Flush` loaded a snapshot, quieted it, and cleared the whole word on the coop's leading thread, on every rank, which could drop a concurrent mark whose doorbell had not yet fired and then skip the quiet that mark needed. Clearing after the drain also keeps the invariant that a clear bit means a drained queue, so a second coop flushing the same context cannot return from `Flush` while the first is still inside `quiet()`. - #5708 (merged 2026-10-29, and its merge commit is not in `ded49fd`), moved from the `RCCL 2.30.7 for ROCm 10.1.0` section: > * `hipcc` is no longer accepted as the CMake CXX compiler for RCCL builds. Use `amdclang++` directly (`-DCMAKE_CXX_COMPILER=amdclang++` or set `ROCM_PATH`). `hipcc` was a wrapper around `amdclang++` that is being phased out across the ROCm stack. Only the CHANGELOG is changed. Version macros are left to the component owners. JIRA ID : AIROCDOC-4467 --------- Co-authored-by: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Follow-up to #12948 (
RCCL_IB_UDMA_POLICY), addressing the review comments that were left after it merged:documentation for the new variable, a guard on the policy enum range, one shared validity check, a consistency
check in AINIC QP recovery, and the QP-count requirement of
SplitConnection.Technical Details
One commit per review item:
ncclIbCastUDMAPolicy(common_cast.h):Lastis nowCount - 1, withstatic_assertsthat
First == ChannelAffinity == 0andLast == SplitConnection. Adding a policy without updating the assertsfails to compile. Nothing
switches over the enum, soCountcannot trigger-Wswitch.ncclIbCastUDMAPolicyIsValid()incommon_cast.h, used by bothinit.cc(WARN +fallback) and
ncclIbCastGetUDMAPolicy()inconnect.cc. No behavior change.devIndexcheck (p2p_resiliency_recovery_ainic.cc): the device index fromIbCastQpStripemustmatch
localQp->devIndex; otherwise WARN and returnncclInternalError. Connect setsdevIndexfrom the samestripe, so this cannot fire today; it catches the striping assumption behind
qpIdxInDevbreaking later.SplitConnectionQP-count note (comments only): the defaultNCCL_IB_QPS_PER_CONNECTIONof 2 alreadysplits a connection. With one QP per device per connection (
NCCL_IB_QPS_PER_CONNECTION=1, orRCCL_IB_QPS_PER_P2P=1for P2P)qpIdxis always 0 and the policy reduces to channel affinity by channel-idpopcount parity. RMA comms also use one QP.
RCCL_IB_UDMA_POLICYentry indocs/api-reference/env-variables.rstand a CHANGELOG "Added" bullet.Policy 0 pins each channel's send QPs, and separately its receive QPs, to one pipeline per device. Policy 1
round-robins P2P QPs on one per-device counter. Policy 3's default of two QPs per connection is enough;
splitting is lost with one QP per connection and on RMA comms.
Default behavior (
RCCL_IB_UDMA_POLICYunset = 0) is unchanged. A per-direction round-robin counter for policy 1was tried and dropped: it did not improve performance.
Issue Tracking
JIRA ID: ROCM-31118
Test Plan
connect.cc,init.cc,p2p_resiliency_recovery_ainic.cc, andthrough them
common_cast.h) on developdf181d33333.Countwithout updating the asserts fails the build.rccl_check_default_behavior.py --base-ref <develop>: passes.Test Result
For reference, #12948's validation on TW (MI355X + AINIC, 8 reps, interleaved arms, 144/144 runs clean per
stage) ran with the default 2 QPs per connection, where this PR does not change the assignment. Policy unset was
flat vs develop on alltoall/all_reduce/all_gather at 2N and 4N. alltoall Avg busbw: 2N +9.3% (policy 1, 8/8),
4N +6.0% (policy 1, 8/8); policies 2 and 3 +3.6% at 4N.
Submission Checklist