Repository navigation
fix(rccl): build the RMSNorm kernel fusion examples on ROCm - #13815
Merged
i-kosarev merged 6 commits intoOct 30, 2026
Merged
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
amd-arozanov
requested changes
Oct 21, 2026
amd-arozanov
left a comment
amd-arozanov
left a comment
Contributor
There was a problem hiding this comment.
- Regarding coop.h fix. I think you can fix it something like that:
#if defined(ROCM_VERSION)
#define NCCL_COOP_HAS_WARP_SYNC (ROCM_VERSION >= 70000)
#elif defined(__HIP_PLATFORM_AMD__)
#include <hip/hip_version.h>
#define NCCL_COOP_HAS_WARP_SYNC (HIP_VERSION_MAJOR >= 7)
#else
#define NCCL_COOP_HAS_WARP_SYNC 0
#endif
...
#if NCCL_COOP_HAS_WARP_SYNC
if (nThreadsPow2 > 1) __syncwarp(laneMask());
#else
__syncthreads();
#endif
- Please also add 07_kernel_fusion to the categories table in examples/README.md and change "build all 01..05 categories". examples/Makefile:11 already builds 06 and 07, and after this PR the top-level make succeeds for 07.
Contributor
Author
|
@amd-arozanov thanks, both done in 2081c68.
|
mkuznet1
marked this pull request as ready for review
October 21, 2026 12:02
![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
2 more finding(s), not posted inline
projects/rccl/src/include/nccl_device/coop.h:19(Low) No test compiles coop.h without ROCM_VERSION, so reverting the application-side NCCL_COOP_HAS_WARP_SYNC branch fails nothingprojects/rccl/src/algorithms/rccl_ep/CMakeLists.txt:56(Nit) rccl_ep/CMakeLists.txt:55-58 still says coop.h falls back to __syncthreads() when ROCM_VERSION is undefined, which this PR changesprojects/rccl/src/algorithms/rccl_ep/CMakeLists.txt:55-58(outside the diff) says coop.h falls back to__syncthreads()whenROCM_VERSIONis undefined and citescoop.h:124 and :154. With this PR coop.h falls back toHIP_VERSION_MAJORinstead. The-DROCM_VERSIONthere is still needed forhip_compat.h:63, so only the coop.h part of the comment needs updating.
Bottom line: Looks approval ready.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
Author
|
ARGUS follow-up in 5939b83. |
Contributor
|
@cppccppccppc @sunguobao PTAL |
amd-arozanov
approved these changes
Oct 21, 2026
garzaran
approved these changes
Oct 22, 2026
i-kosarev
force-pushed
the
users/mkuznet1/aicomrccl-1370-rmsnorm-examples
branch
from
October 28, 2026 05:44
48d2825 to
abec5d1
Compare
cmcknigh
approved these changes
Oct 28, 2026
cmcknigh
left a comment
cmcknigh
left a comment
Contributor
There was a problem hiding this comment.
percheckin change seems fine.
Contributor
pr-rocm-systems-13815-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. |
Contributor
pr-rocm-systems-13815-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-13815-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. |
Contributor
pr-rocm-systems-13815-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. |
Contributor
pr-rocm-systems-13815-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
i-kosarev
force-pushed
the
users/mkuznet1/aicomrccl-1370-rmsnorm-examples
branch
from
October 29, 2026 11:06
abec5d1 to
0a529db
Compare
Contributor
pr-rocm-systems-13815-oci-rccl-pr-buildBuilds
|
Contributor
pr-rocm-systems-13815-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
i-kosarev
force-pushed
the
users/mkuznet1/aicomrccl-1370-rmsnorm-examples
branch
from
October 29, 2026 13:14
0a529db to
808e64a
Compare
Contributor
pr-rocm-systems-13815-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. pr-rocm-systems-13815-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-13815-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
The 07_kernel_fusion examples came in with NCCL 2.30.7 but their Makefiles still included the upstream NVIDIA build rules, which resolve outside the repo and compile .cu with nvcc. Point them at rccl_examples.mk, which hipifies the unmodified upstream main.cu at build time, as 06_device_api already does. rmsnorm_utils.cuh includes <cuda_runtime.h>, so add a rule that hipifies .cuh headers into HIPIFY_DIR as well; the hipified main.cu finds that copy first through its quoted include. Ignore the built example binaries. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
nccl_device/coop.h picked __syncwarp from ROCM_VERSION, which only RCCL's own build defines. Application kernels including the installed headers fell back to __syncthreads() in the thread and warp coops; the GIN barrier calls those from a subset of threads, so the block's waves went out of step and the 07_kernel_fusion GIN/hybrid examples returned wrong results. Fall back to HIP_VERSION_MAJOR when ROCM_VERSION is not defined. Also list 07_kernel_fusion in the examples README and drop the CUDA requirements from the RMSNorm example Makefile help. Co-authored-by: Cursor <cursoragent@cursor.com>
Add a device test that compiles nccl_device/coop.h the way an application does, with RCCL's ROCM_VERSION undefined. A static_assert pins the gate to the HIP version, and three kernels run ncclCoopThread, ncclCoopWarp, and ncclCoopLanes sync in one wave only; if any of them is a block barrier, the block reduction reads a slot before that wave publishes it. Shorten the hipify .cuh rule and coop.h comments, and update the rccl_ep comment that still described the old __syncthreads() fallback. Co-authored-by: Cursor <cursoragent@cursor.com>
clock64() is the 20-bit SHADER_CYCLES counter on gfx11, so the delay loop could wrap and never exit. wall_clock64() is 64-bit on every arch. Co-authored-by: Cursor <cursoragent@cursor.com>
The precheckin config lists rccl-UnitTestsFixtures filters explicitly, so the new coop test was built but never run. Co-authored-by: Cursor <cursoragent@cursor.com>
i-kosarev
force-pushed
the
users/mkuznet1/aicomrccl-1370-rmsnorm-examples
branch
from
October 30, 2026 11:10
808e64a to
57d21a3
Compare
Contributor
pr-rocm-systems-13815-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-13815-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
TheRock Submodule Bump ActivityNewest first
|
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>
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
NCCL 2.30.7 added the RMSNorm kernel fusion examples (
examples/07_kernel_fusion: LSA, Multimem, GIN, Hybrid LSA/GIN). They came into RCCL with the sync but still included the upstreammakefiles/{common,examples}.mk(outside the repo) and built withnvcc, so none of them could be built on ROCm. This PR wires them up the same way #11579 did for06_device_api.Validating them also exposed a defect in the public device API headers that affects every application kernel using the GIN device API, not only these examples. This PR fixes it as well (see Device API fix below).
Technical Details
c/Makefiles includeexamples/makefiles/rccl_examples.mk. The.cusources are hipified at build time withhipify-perlinto.hipify/and stay byte-identical to NCCL, so future syncs apply cleanly.rccl_examples.mk: new$(HIPIFY_DIR)/%.cuhrule. The sharedinclude/rmsnorm_utils.cuhincludes<cuda_runtime.h>, so it is hipified too. The quoted#includein the hipifiedmain.cusearches.hipify/first and picks up the HIP copy.RCCL_HOME,ROCM_PATH,GPU_TARGETS), and the Makefilehelptargets drop the CUDA requirements, as in06_device_api.examples/README.mdlists07_kernel_fusion, and its build line covers 01..07, matchingexamples/Makefile.examples/.gitignore.Device API fix:
nccl_device/coop.hdepended on a macro application code never seescoop.hselected__syncwarpwith#if ROCM_VERSION >= 70000and otherwise fell back to__syncthreads().ROCM_VERSIONis defined only by RCCL's own CMake (CMakeLists.txt:393). Applications that include the installednccl_deviceheaders do not have it, soncclCoopTile::sync()(includingncclCoopThread) andncclCoopLanes::sync()became block-wide barriers.The GIN barrier (
impl/gin_barrier__funcs.h:98-115) callssignal/waitSignalwithncclCoopThread()from the threads withpeer < nRanksonly, i.e. part of wave 0. Wave 0 then executed extras_barriers and ran out of phase with the other waves:03_rmsnorm_ginand04_rmsnorm_hybridreturned wrong results: the block reduction inblockRMSNormsaw only wave 0's partial sums (1/4 of the sum of squares with 256 threads, so every token came out about 2x too large).06_device_api/02_alltoall_ginstill passed because it checks the data on the host after the kernel. It does not consume the received data inside the kernel.coop.hnow usesNCCL_COOP_HAS_WARP_SYNC:ROCM_VERSION >= 70000inside RCCL's build (unchanged),HIP_VERSION_MAJOR >= 7for application code. RCCL's own kernels compile exactly as before. CHANGELOG: 2.32.3, Resolved issues.Issue Tracking
JIRA ID: AICOMRCCL-1370.
Test Plan
With RCCL built and installed from this branch, from
examples/(top-levelmake) and from07_kernel_fusion/06_device_api(06as a control):The examples exit 0 even when the check fails, so the result is read from the output (
SUCCESS:/FAILED:/ERROR:).Test Result
Build: the four RMSNorm examples build through
hipify-perl(sources and the shared header) with no warnings, from07_kernel_fusionand from the top-levelmake. The top-levelmakestill stops in 01..05: their leafc/Makefiles include the upstreammakefiles/examples.mk, which RCCL does not carry. That predates this PR and is not changed here.2x 8x MI350X (gfx950, ROCm 7.0.2, RoCE
bnxt_re), bare metal, no extra compiler flags:coop.hfix (f21b1b92cc)2081c68d56)communicator does not support MultimemMultimem is not available on AMD GPUs (no NVLS; the multimem PTX path is compiled only for
__CUDA_ARCH__ >= 900), so02_rmsnorm_multimemstops with that message by design.8x MI355X (gfx950), single node, containerized: 01 LSA SUCCESS on 8 and 2 GPUs. The GIN examples could not be validated there: RDMA inside that container is limited, and the
06GIN control also hangs with 8 ranks.