Repository navigation
fix(rccl/gin): drop Anvil SDMA backend on multi-node communicators - #14216
Merged
dlamd1dai merged 9 commits intoOct 31, 2026
Merged
Conversation
With NCCL_GIN_TYPE unset, Anvil SDMA sorted ahead of PROXY and became the default backend on every communicator. Its factory addresses peers by local GPU device id, so on more than one node a remote rank maps to a local queue. Remove it in ncclGinSetDefaultBackend when nNodes > 1, before connect, so PROXY is the default. Single-node ordering is unchanged. Add GinMPIDeviceTests.MixedBackends_TwoNodeUnsetType: 2 ranks on 2 nodes, checks props.ginType/ginSupport, rejects an explicit Anvil DevComm, and runs put/signal on a default and an explicit PROXY DevComm. AICOMRCCL-3417 Co-authored-by: Cursor <cursoragent@cursor.com>
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Contributor
pr-rocm-systems-14216-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-14216-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
![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/test/transport/GinDeviceMPITests.cpp:151(Low) crossNodeReason() in mixedBackendTwoNodeSkipReason duplicates the min_nodes=2 check at line 7897; drop it.
Bottom line: Worth settling before approval: dropping Anvil can leave zero GIN backends while globalGinSupport stays FULL/RAIL (gin.cc:461), and the new MPI test never runs in CI and passes without the fix when Anvil is not built.
Comment @argus review to re-run, or @argus re-review after pushing.
On gfx1250, MNNVL (UALoE) makes a pod one node, and Anvil SDMA reaches every GPU in it, so comm->nNodes > 1 is the right cutoff. Two hosts in one pod give nNodes == 1; skip the case there instead of failing its nNodes assertion, and note SDMA's reach next to the guard. AICOMRCCL-3417 Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
pr-rocm-systems-14216-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-14216-ruby64-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
This was referenced Oct 25, 2026
Drop Anvil when topo reports multiple physical hosts so multi-host MNNVL cannot keep a device-id local backend, and only set globalGinSupport when SetDefaultBackend left at least one backend. Cover both branches in the gin-plugin host suite and skip the MPI case unless Anvil is actually built. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
@argus re-review |
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: the new Anvil host tests never run in CI (gin-plugin-init-test.cc:254), and no test covers the new supported gate at init.cc:2869.
Comment @argus review to re-run, or @argus re-review after pushing.
1 task done
Contributor
pr-rocm-systems-14216-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. |
…ver the empty-backend gate NCCL_GIN_TYPE_ANVIL_SDMA is a plain enum, so the multi-host drop in ncclGinSetDefaultBackend and its three GinFinalizeTest cases no longer need ENABLE_ROCSHMEM_GIN; precheckin's CPU-only build now runs them. The FULL or RAIL choice after ncclGinSetDefaultBackend moves into ginConnectionType, tested next to ginGateAllows, so removing the ginState.supported check fails a test. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: gin.cc:455 and GinDeviceMPITests.cpp:7914 read comm->topo->nHosts without including graph/topo.h, so librccl and the MPI test binary hit an incomplete-type error.
Comment @argus review to re-run, or @argus re-review after pushing.
aduforest
reviewed
Oct 28, 2026
ncclTopoSystem is only forward-declared from graph.h, so reading nHosts in gin.cc and the mixed-backend MPI test failed as an incomplete type. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
@argus re-review |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-14216-ruby64-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
aduforest
approved these changes
Oct 28, 2026
Contributor
Author
|
@ROCm/rccl-ci requesting an override of the required RCCL PR gate check. It fails only because of
This PR is approved. It has to merge before #13429. |
This was referenced Oct 28, 2026
Contributor
Author
|
@abouteiller could you take a look when you have a moment? This PR needs one rccl-leads approval to merge. It has to land before #13429 (which you approved): with |
abouteiller
approved these changes
Oct 29, 2026
…omrccl-3417-mixed-gin-backends Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/rccl/test/transport/GinDeviceMPITests.cpp
Contributor
Author
|
Thanks @abouteiller and @aduforest for the reviews. I merged @ROCm/rccl-ci this restarts CI. Once Ruby smoke finishes on 4c0b63b, could you apply the RCCL PR gate override to the new commit if it fails on the same |
Contributor
pr-rocm-systems-14216-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 12 tests skipped. |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 24 tests skipped. |
Contributor
pr-rocm-systems-14216-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 12 tests skipped. |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 24 tests skipped. |
Contributor
pr-rocm-systems-14216-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
|
PR #14216 – fix(rccl/gin): drop Anvil SDMA backend on multi-node communicators: passed at head. Earlier red runs are superseded. Failing check: none on head 7892bbe. Run 41012466415 — Ruby smoke, OCI smoke, and the RCCL PR gate all succeeded. Test: none. Ruby session f6f4a6ec is 275 passed and 15 skipped. OCI session 09eb8bac is 254 passed and 15 skipped. The build sessions are all skipped, which is what --build-only records. Error: none in the head run. Evidence: SHAs 6b27a3a and 4c0b63b failed both clusters and have build sessions only. Before that, the same 4-node regression perf rows (alltoall_perf, reduce_scatter_perf, all_gather_perf, all_reduce_perf) failed on the PR stacks and on rocm-10.1-ga_rccl-develop-perf (OCI 451b56cc on 2026-10-24; Ruby 6d6cee73 on 2026-10-25 and d3bb21d4 on 2026-10-28). Those rows passed on the later sessions, including the head. Existing ticket: Jira not checked. The head run has no failing test. Next step: nothing for coco on this SHA. |
…omrccl-3417-mixed-gin-backends
Contributor
pr-rocm-systems-14216-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14216-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
4 of 6 tasks
dlamd1dai
added a commit
that referenced
this pull request
Oct 31, 2026
This PR is rebased onto `develop` (`e7cb4b477a`, head `126528dfdf`; only the CHANGELOG context changed since the tested head `7237ae7b96`) after #12751, and the check lives in that PR's `ginAnvilRegisterSignalSpan`. **Merge #14216 first:** with `NCCL_GIN_TYPE` unset on a multi-node communicator, Anvil can still be the selected backend until #14216 drops it there, and this check would then fail `ncclDevCommCreate` for Device API users instead of the communicator falling back to the GIN proxy. #12841 and #12742 from the CPX stack have merged; this PR does not depend on #13235 or #13237 and can merge in any order with them. The defect was found while fixing AICOMRCCL-2387 and is follow-up 4 on #12742. The CPX hardware runs below used the CPX stack so that Anvil could initialize on partitions. ## Motivation On MI300X CPX with 8 ranks and `NCCL_LSA_TEAM_SIZE=4`, the Anvil SDMA GIN backend (`NCCL_GIN_TYPE=7`) publishes signal addresses that are not mapped on the device. At init it only warns: ``` signal nRanks=8 != devr->lsaSize=4 ctx->rank=4 != devr->lsaSelf=0 ``` It then skips the connectivity check and keeps going. The signal binding path (`ginAnvilBuildSignalPeerAddrs`, called from `ginAnvilRegisterSignalSpan` and `ginAnvilBindContextSignals`) builds the peer signal table as `lsaSelf + (pe - ginRank) * bigSize` for every GIN peer. The device indexes this table by GIN peer (`remoteSignalAddr`, `resolveRemotePeerVa`), but each entry is a slot in this rank's LSA window. That is only correct for peers in the same LSA team, where GIN position and LSA slot agree. On rank 0, entries for GIN peers 4–7 point past the end of the window. On rank 4, entries for peers 0–3 point before its start. The ticket suggests rebasing the table on the LSA origin and extent. That would move the in-team peers' entries off their real slots, so this PR does not do that. Writing 0 for out-of-team peers is not safe either: `signalPeer` returns silently on a null address, so the wait would hang, and `ginAnvilResolvePeerVa` turns a null base into `0 + offset`. The same problem affects any `NCCL_GIN_CONNECTION_RAIL` GIN team with more than one rank. Its peers are on other hosts and never have a slot in the local window. This PR makes these configurations fail at registration with a clear error. The symmetric kernels, which request GIN internally and whose GIN peers are always rail peers outside the LSA team, now skip GIN on Anvil, so applications that never ask for GIN are not affected. ## Technical Details - `src/gin/gin_plugin_anvil_sdma.cc`: - New helper `ginAnvilFirstSignalPeerOutsideLsa(comm, nRanks, ginRank)`. For each GIN peer, it computes the LSA slot that the table entry points at. It accepts the peer only when that slot is inside `[0, lsaSize)` and `lsaRankList[slot]` is the peer's world rank. The GIN-to-world mapping uses the same stride as `teamRankToGinRank`: `contiguousRanksPerHost` under RAIL, otherwise 1. - `ginAnvilRegisterSignalSpan` calls the helper before it builds the peer table or registers the span in the IPC table. It runs once per resource-window bind, so it covers every logical context of every pending GIN context. If a peer fails, it logs a `WARN` naming the peer, GIN rank and size, `lsaSelf` and `lsaSize`. The warning tells the user to set `NCCL_LSA_TEAM_SIZE` to cover the GIN team or select another `NCCL_GIN_TYPE`. Registration then returns `ncclInvalidUsage`, which `ncclGinAnvilBindResourceWindowSignals` propagates to the caller that created the device communicator (see the next list). - The two old warnings are removed. Communicators whose GIN team matches their LSA team take the same path as before. - `test/gin/GinAnvilPlugin_test.cpp`: - New helper `expectSignalsUnbound`. It checks that a failed registration left `signals` and `signal_remote_addrs` null in every logical context of the handle, added nothing to the IPC table, and ran no connectivity check. - `ConnCheck_SkippedWhenLsaSizeMismatch` is now `BindSignals_GinPeerPastLsaWindowFailsRegistration`. - New test `BindSignals_GinPeerBeforeLsaWindowFailsRegistration`: the CPX case seen from the second LSA team (rank 2 of 4, `lsaSize` 2). - `ConnCheck_SkippedOnRailConnection` is now `BindSignals_RailGinTeamFailsRegistration`. In its layout, the rail peer's slot is inside the window but holds the wrong world rank, so the test only passes if the RAIL stride is applied. - `ConnCheck_UsesLsaTeamForBootstrap` used a reversed `lsaRanks` list that no real communicator produces. It now uses rank 1 with `lsaSelf` 1 and checks that the bootstrap uses the intra-node rank and peer list. - `connectColl` now takes the GIN rank. - `BindSignals_SpanBaseProbeSizedForWorldRanks` (added by #12751) used a two-rank RAIL GIN team, which now fails registration. It now uses one host of 8 ranks, where the rail team is this rank alone, so the world (8 ranks) is still wider than the GIN team (1 rank) and the probe buffer sizing is still tested. - `src/sym_kernels.cc`, `src/include/sym_kernels.h`: - New `ncclSymkState::ginBackendUnsupported`. `ncclSymkInitOnce` sets it when the symmetric kernels need GIN (LSA team smaller than the communicator) and `ncclGetRailedGinType` reports `NCCL_GIN_TYPE_ANVIL_SDMA`. In that case it logs an `INFO` line and does not add the GIN requirement to its internal device communicator, so Anvil signals are never bound during init. - `ncclSymkMask` treats GIN as unavailable when the flag is set. Because the mask keeps only GIN kernels when GIN is needed, this disables the symmetric kernels for that communicator, the same as `NCCL_SYM_GIN_KERNELS_ENABLE=0`; collectives use the regular kernels. - `ncclGinSetDefaultBackend` runs before `ncclSymkInitOnce` in `init.cc`, so the backend is known when the flag is computed. - `test/host/sym-kernels-test.cc`: replaces `ncclGetRailedGinType` with a hook. `SymKernelMicrotest.GinBackendUnsupported_OnlyForAnvilSdma` checks that only Anvil sets the flag (Anvil, proxy, none). `SymKernelMaskTest.NeedGinTrue_GinBackendUnsupported_EverythingCleared` checks that the mask is empty when GIN is needed and the flag is set. - `CHANGELOG.md`: added a "Resolved issues" entry. It notes that `ncclDevCommCreate` calls requesting GIN in these configurations used to succeed and now fail, and that the symmetric kernels no longer use GIN with Anvil, so communicator init and collectives are unaffected. Behavior changes reviewers should know about: - **Narrowed LSA team, including CPX across OAMs:** with `NCCL_GIN_TYPE=7` and `NCCL_LSA_TEAM_SIZE` smaller than the communicator, `ncclDevCommCreate` calls that request GIN used to succeed and now fail with `ncclInvalidUsage` and the new warning. `ncclCommInitRank` and collectives are unaffected: the symmetric kernels skip GIN on Anvil (see `sym_kernels.cc` above), so for these communicators they are disabled instead of failing init. On hardware, a plain `all_gather_perf` in this setup passes on this PR (see Test Result). - **RAIL with more than one GIN rank:** Device API GIN requests always fail on Anvil now. The symmetric kernels never use Anvil for it. - **Multi-backend runs:** with `NCCL_GIN_TYPE` unset, Anvil can be one of several active backends and ranks ahead of the GIN proxy. `ncclGinDevCommSetup` falls back between backends, but signal binding happens after that choice and has no fallback. #14216 drops Anvil on multi-node communicators so the proxy is chosen there; merge it first. On a single node with a narrowed LSA team, a Device API communicator that selected Anvil now fails instead of signaling through unmapped addresses. This has not been run multi-node. - **Connectivity check:** a communicator that passes registration reaches the skip branches for "GIN team size differs from LSA team size" and for RAIL only when its GIN team is smaller than its LSA team, or is a single-rank rail team. They are left in place. Related: - #14216 (merge first) drops Anvil when `comm->topo->nHosts > 1`. - #12751 (merged) rewrote this function as `ginAnvilBuildSignalPeerAddrs` / `ginAnvilRegisterSignalSpan`. The rebase carries the check into `ginAnvilRegisterSignalSpan` and removes the old warnings there; the second commit adapts the tests to per-context signal spans. The branch is now three commits; an attempt to refuse out-of-team peers in `createContext`, and its revert, were dropped. - The buffer path's IPC-table fallback in `regMrSym` (`RegisterVmm` with `lsaSelf` / `lsaSize`) has the same index-space mismatch. The stride path in `resolveRemotePeerVa` usually bypasses it, and this PR does not change it. ## Issue Tracking JIRA ID: AICOMRCCL-2706 ## Test Plan - `rccl-UnitTestsGinAnvilPlugin --gtest_filter='*BindSignals*:*ConnCheck*'`, then the full `rccl-UnitTestsGinAnvilPlugin` suite. Requires `ENABLE_ROCSHMEM_GIN` and ROCm 6.4 or later. - `rccl-UnitTestsMicroSymKernels` (host microtests). - rccl-tests `alltoall_perf` and `all_gather_perf` with `NCCL_GIN_TYPE=7`, on CPX ranks split across two OAMs (MI355X and MI300X) and on SPX (MI300X), each on this PR's library and on a baseline library with the plugin file from `develop`: - LSA team narrowed below the communicator: expect init to fail with the new warning instead of the old "nRanks != lsaSize" warnings. - LSA team unset: expect init and the connectivity check to pass, and `-D 3` runs to complete with the data check passing. - Plain collectives (no `-D`) with a narrowed LSA team: expect them to pass, with the symmetric-kernel skip logged on every rank. ## Test Result On MI355X with Linux 6.8, Anvil passes `alltoall_perf -D 3` and `all_gather_perf -D 3` data checks when the GIN and LSA teams match, in SPX with 8 ranks and in CPX with 4. When the LSA team is narrowed, this PR fails `ncclDevCommCreate` on every rank with the new warning, and plain collectives pass. On the same layout, the baseline skips the connectivity check and then hangs or fails. The MI300X host repeats the failure-path results. ### With the symmetric-kernel skip (`develop` `daf7ef33e4`, head `7237ae7b96`) Same MI355X host and build as the next section, SPX, 8 ranks, this PR's library only. Logs are under `~/pr13429/logs-port2`. - Full `rccl-UnitTestsGinAnvilPlugin`: 48 of 48 passed. - `rccl-UnitTestsMicroSymKernels`: 69 of 69 passed, including the two new tests. | Run | GIN team vs. LSA team | Result | Symmetric GIN skip logged | |---|---|---|---| | `alltoall_perf -D 3`, `NCCL_LSA_TEAM_SIZE=8` | 8 vs. 8 | Signals bound, conn-check OK on all 8 ranks, 0 wrong values | 0 ranks | | `all_gather_perf -D 3`, `NCCL_LSA_TEAM_SIZE=8` | 8 vs. 8 | Signals bound, conn-check OK on all 8 ranks, 0 wrong values | 0 ranks | | `alltoall_perf -D 3`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | Init passes; the `ncclDevCommCreate` group fails (`common.cu:2705`, `invalid usage`) with the new warning on all 8 ranks | 8 ranks | | `all_gather_perf -D 3`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | rccl-tests refuses the kernel before creating a device communicator ("GinHybridAllGatherKernel requires CUDA P2P connectivity across all ranks"), rc 6; the baseline library prints the same refusal | 8 ranks | | Plain `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | **Passes**, 0 wrong values (failed `ncclCommInitRank` at the previous head) | 8 ranks | | Plain `all_gather_perf -R 2`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | Passes, 0 wrong values | 8 ranks | | Plain `alltoall_perf -R 2`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | Passes, 0 wrong values | 8 ranks | Example lines: ``` NCCL INFO Symmetric GIN kernels disabled: the Anvil SDMA GIN backend cannot reach rail peers outside the LSA team (lsaSize=4, nRanks=8) NCCL WARN GIN anvil-sdma: GIN peer 4 is outside the LSA team of rank 2 (ginRank=2, ginRanks=8, lsaSelf=2, lsaSize=4). Anvil SDMA can only signal LSA team peers; set NCCL_LSA_TEAM_SIZE to cover the GIN team or select another NCCL_GIN_TYPE. ``` The sections below are from earlier heads, before the symmetric-kernel skip. Where they say "Fails in `ncclCommInitRank`", this head passes instead. ### After the rebase onto `develop` (`1ba7048298`, head `2cb35ac234`) Built on the MI355X host below (gfx950, ROCm 10.1.0, Linux 6.8.0-101) with `install.sh --rocshmem-gin`, then `-DBUILD_TESTS=ON -DGIN_ANVIL_UNIT_TESTS=ON`. The baseline library uses `gin_plugin_anvil_sdma.cc` from `develop` (`1ba7048298`), which includes #12751. Logs are under `~/pr13429/logs-port` on that host. - `rccl-UnitTestsGinAnvilPlugin --gtest_filter='*BindSignals*:*ConnCheck*'`: 30 of 30 passed. - Full `rccl-UnitTestsGinAnvilPlugin`: 48 of 48 passed. **SPX, 8 ranks, one per GPU:** | Run | GIN team vs. LSA team | This PR | Baseline | |---|---|---|---| | `alltoall_perf`, `NCCL_LSA_TEAM_SIZE=8` | 8 vs. 8 | Signals bound (4 logical contexts per rank), conn-check **OK** on all 8 ranks, 0 wrong values | Same as this PR | | `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=8` | 8 vs. 8 | Signals bound, conn-check OK on all 8 ranks, 0 wrong values | Same as this PR | | `alltoall_perf`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | Init fails with `invalid usage` on all 8 ranks, new warning | Old warnings and conn-check skipped on all ranks, then **hangs** (killed at 300 s) | | `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | Init fails, new warning on all 8 ranks | Old warnings and conn-check skipped, then the test fails (rc 6) | | Plain `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | Fails in `ncclCommInitRank`, new warning | Passes, 0 wrong values, old warnings printed | | Plain `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=4`, `NCCL_SYM_GIN_KERNELS_ENABLE=0` | 8 vs. 4 | Passes, 0 wrong values | Passes | | `alltoall_perf`, `NCCL_LSA_TEAM_SIZE=4`, `NCCL_SYM_GIN_KERNELS_ENABLE=0` | 8 vs. 4 | Fails in the `ncclDevCommCreate` group (`common.cu:2705`), new warning | (not run) | The CPX cases were not re-run after the rebase, because the host was in SPX and changing its partition mode would disturb other users. The tables below are from the pre-rebase head. ### Before the rebase **Unit tests pass** (MI300X host below, ROCm 7.2.0, built at `9e444f3ab5` with `install.sh --rocshmem-gin`, then `-DBUILD_TESTS=ON -DGIN_ANVIL_UNIT_TESTS=ON`): - `rccl-UnitTestsGinAnvilPlugin --gtest_filter='*BindSignals*:*ConnCheck*'`: 23 of 23 passed, including the three new `BindSignals_*FailsRegistration` tests. - Full `rccl-UnitTestsGinAnvilPlugin`: 37 of 37 passed. Each hardware case ran twice: on this PR's library, and on a baseline library built from the same tree with `gin_plugin_anvil_sdma.cc` from `develop` (`2a4afbe5c4`). Common settings were `NCCL_GIN_TYPE=7`, `NCCL_CUMEM_ENABLE=1` and `-R 2 -D 3`, unless a row says "plain" (no `-D`). Runs were killed at 300 s. ### MI355X: `smci355-ccs-aus-m03-17` 8× MI355X (gfx950), NPS1, Linux 6.8.0-101, amdgpu 7.1-dev, ROCm 10.1.0, built at `8a97123fff`. No kernel-gate bypass was needed. Logs are under `~/pr13429/logs` on that host. **SPX, 8 ranks, one per GPU** (this PR's tree only): | Run | GIN team vs. LSA team | This PR | Baseline | |---|---|---|---| | `alltoall_perf`, `NCCL_LSA_TEAM_SIZE=8` | 8 vs. 8 | Signals bound (`lsaSize=8`), conn-check **OK** on all 8 ranks, 0 wrong values | Same as this PR | | `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=8` | 8 vs. 8 | Signals bound, conn-check OK on all 8 ranks, 0 wrong values | Same as this PR | | `alltoall_perf`, `NCCL_LSA_TEAM_SIZE=4` | 8 vs. 4 | Init fails with `invalid usage` on all 8 ranks, new warning | (not run) | **CPX, 4 ranks** (this PR plus the CPX stack: #12841, #13235, #13237 and #12742 at their 2026-10-14 heads, when the even-engine change was still in #13237). Ranks 0–1 are on two partitions of one OAM and ranks 2–3 on two partitions of another (`HIP_VISIBLE_DEVICES=0,1,8,9`, `NCCL_MULTI_RANK_GPU_ENABLE=1`): | Run | GIN team vs. LSA team | This PR | Baseline | |---|---|---|---| | `alltoall_perf`, LSA team unset | 4 vs. 4 | Signals bound, conn-check **OK** on all 4 ranks, 0 wrong values | Same as this PR | | `all_gather_perf`, LSA team unset | 4 vs. 4 | Signals bound, conn-check OK on all 4 ranks, 0 wrong values | Same as this PR | | `alltoall_perf`, `NCCL_LSA_TEAM_SIZE=2` | 4 vs. 2 | Init fails with `invalid usage` on all 4 ranks, new warning | Old warnings and conn-check skipped on all ranks, then **hangs** (killed at 300 s) | | `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=2` | 4 vs. 2 | Init fails, new warning on all 4 ranks | Old warnings and conn-check skipped, then the test fails (rc 6) | | Plain `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=2` | 4 vs. 2 | Fails in `ncclCommInitRank`, new warning | Passes, 0 wrong values, old warnings printed | | Plain `all_gather_perf`, `NCCL_LSA_TEAM_SIZE=2`, `NCCL_SYM_GIN_KERNELS_ENABLE=0` | 4 vs. 2 | Passes, 0 wrong values | Passes | | `alltoall_perf`, `NCCL_LSA_TEAM_SIZE=2`, `NCCL_SYM_GIN_KERNELS_ENABLE=0` | 4 vs. 2 | Fails in the `ncclDevCommCreate` group (`common.cu:2705`), new warning | (not run) | Example warning (rank 2, on the second OAM): ``` GIN anvil-sdma: GIN peer 0 is outside the LSA team of rank 2 (ginRank=2, ginRanks=4, lsaSelf=0, lsaSize=2). Anvil SDMA can only signal LSA team peers; set NCCL_LSA_TEAM_SIZE to cover the GIN team or select another NCCL_GIN_TYPE. ``` The baseline prints the warnings from the ticket on the same run, for example `signal nRanks=4 != devr->lsaSize=2 (rank 0)` and `ctx->rank=2 != devr->lsaSelf=0 (signal peer indexing may be wrong)`. ### MI300X: `banff-cyxtera-cx56-5` 8× MI300X (gfx942), NPS1, Linux 5.15, amdgpu 6.16.13, ROCm 10.1. Logs are under `~/pr13429/logs-r10` on that host. On this host, RCCL turns cuMem off below Linux 6.8, and without cuMem rccl-tests skips every `-D 3` run. These runs therefore add an experiment-only patch that skips that kernel gate (`RCCL_EXP_SKIP_KERNEL_GATE=1`). On this kernel, cuMem peer mappings are broken, so the matched-team runs fail the connectivity check on both libraries. The failure path still shows, because the new check runs before any peer is signaled. `NCCL_RMA_DISABLE=1` was used whenever the LSA team was narrowed, because the RMA proxy's memory registration fails on this host before GIN is reached. | Run | GIN team vs. LSA team | This PR | Baseline | |---|---|---|---| | SPX, 8 ranks, `NCCL_LSA_TEAM_SIZE=4`, `alltoall_perf` | 8 vs. 4 | Init fails with `invalid usage` on all 8 ranks with the new warning | Old warnings and conn-check skipped on all ranks, then **hangs** (killed at 300 s) | | SPX, same, `all_gather_perf` | 8 vs. 4 | Init fails, new warning on all 8 ranks | Old warnings and conn-check skipped, then the test fails (rc 6) | | CPX, 6 ranks, 3 on each of two OAMs, `NCCL_LSA_TEAM_SIZE=3`, `alltoall_perf` | 6 vs. 3 | Init fails, new warning on all 6 ranks | Old warnings and conn-check skipped, then **hangs** (killed at 300 s) | | CPX, same 6 ranks, LSA team unset | 6 vs. 6 | Signals bound, conn-check runs and fails on this kernel's broken cuMem mapping | Same as this PR | | SPX, 8 ranks, LSA team unset | 8 vs. 8 | Signals bound, conn-check runs and fails on the broken mapping | Same as this PR | | SPX, 8 ranks, `NCCL_LSA_TEAM_SIZE=4`, plain `all_gather_perf` | 8 vs. 4 | Fails in `ncclCommInitRank` | Passes, 0 wrong values, old warnings printed | | Same, `NCCL_SYM_GIN_KERNELS_ENABLE=0` | 8 vs. 4 | Passes | Passes | | SPX, 8 ranks, `NCCL_LSA_TEAM_SIZE=4`, `alltoall_perf`, `NCCL_SYM_GIN_KERNELS_ENABLE=0` | 8 vs. 4 | Fails in the `ncclDevCommCreate` group | (not run) | | SPX, 8 ranks, `NCCL_LSA_TEAM_SIZE=4`, plain `all_gather_perf`, `NCCL_GIN_TYPE` unset | 8 vs. 4 | Init fails in the GIN proxy backend (`gin_host_proxy.cc:477`), before Anvil binds | Same failure | ### Notes - The ticket's exact layout (8 CPX ranks, 4 on each OAM) could not be run on either host. On the CPX stack, every rank runs out of SDMA queues while connecting to 7 peers (`SDMA queue budget exhausted: 7 queue(s) taken by this process + 1 requested, limit 8`; on MI355X the driver returns `NO_MEMORY` for the 8th queue). That happens in `connect`, before signal binding, on this PR and on the baseline alike. On MI355X it happens with `NCCL_LSA_TEAM_SIZE=8` too, whether the 8 partitions are on one OAM, split 4+4, or spread one per OAM. The 4-rank (MI355X) and 6-rank (MI300X) layouts stand in for it. The limit of 8 is the even-engine CPX budget, which #12841 now ships; #13237's slot-aware budget raises it to 12, and the 8-rank layout has not been retried with that budget. - The CPX stack needed one local build fix at the time: RCCL's copy list lacked `anvil_engine_map.hpp`. #12841 now includes it. - Not run: multi-node, and the RAIL connection type on hardware. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
that referenced
this pull request
Oct 31, 2026
…omrccl-3418-gda-alongside Resolve conflicts with #14216 (AICOMRCCL-3417): - gin-plugin-init-test.cc: keep the NCCL_GIN_TYPES and backend-order tests and the new multi-host Anvil SDMA tests. - GinDeviceMPITests.cpp: keep the runPutWaitRoundTrip comment that describes the backend check and the rank-1-only asserts, and agree on HasFatalFailure() between the two round trips in AnvilSdma_DroppedOnMultiHost so a rank-1 failure cannot leave rank 0 waiting in the second one. Co-authored-by: Cursor <cursoragent@cursor.com>
Open
1 of 3 tasks
Contributor
TheRock Submodule Bump ActivityNewest first
|
2 of 4 tasks
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
Split from AICOMRCCL-1999 (per-DevComm GIN backend selection, NCCL 2.31.2).
In a build with
ENABLE_ROCSHMEM_GINandNCCL_GIN_TYPEunset, Anvil SDMA and PROXY are both active GIN backends. SDMA has the higher priority, so it becomesbackends[0](the default, and whatprops.ginTypereports) on every communicator, including multi-host ones. The SDMA factory identifies peers by local GPU device id, so a rank on another host with the same device id is mapped to a local SDMA queue.JIRA ID: AICOMRCCL-3417
Technical details
ncclGinSetDefaultBackend(src/plugin/gin.cc) now finalizes and removesNCCL_GIN_TYPE_ANVIL_SDMAwhencomm->topo->nHosts > 1. This covers multi-host MNNVL wherenNodesfolds to 1. Removal happens beforencclGinConnectOnce, so SDMA is never connected across hosts and PROXY becomes the default. The removal is logged at INFO. Single-host backend set and ordering are unchanged.initTransportsRank(src/init.cc) now picksglobalGinSupportthrough a newginConnectionType()helper. It returns NONE whenginState.supportedis false after that filter, so an Anvil-only multi-host communicator reports NONE instead of FULL/RAIL with zero backends.GinMPIDeviceTests.MixedBackends_TwoNodeUnsetType(test/transport/GinDeviceMPITests.cpp). It runs with 2 ranks on 2 hosts andNCCL_GIN_TYPEunset (and skips unless Anvil SDMA is built and probeable), and checks:props.ginType == PROXY,ginSupport[PROXY]set,ginSupport[ANVIL_SDMA]clearginType = ANVIL_SDMADevComm is rejectedENABLE_ROCSHMEM_GIN):GinFinalizeTest.AnvilSdma_*(test/host/gin-plugin-init-test.cc): Anvil kept on one host, dropped with PROXY remaining on two hosts, and Anvil-only clearingsupported.InitMicrotest.GinConnectionType_*(test/host/init-test.cc): NONE when no backend is left, FULL/RAIL otherwise.runPutWaitRoundTriptakes an optionalginType(defaultNCCL_GIN_TYPE_NONE, so existing callers are unchanged).Test result
2 × MI300X (gfx942),
ctr-cx65-mi300x-3+ctr-cx66-mi300x-3, 1 rank per node, CX7 IB, Debug build withENABLE_ROCSHMEM_GIN=ONandENABLE_MPI_TESTS=ON,NCCL_CUMEM_ENABLE=1,NCCL_GIN_TYPEunset:That run was on the AICOMRCCL-3180 branch with an earlier revision of this change, which gated on
nNodes. The current code gates onnHosts, so the INFO line now readssingle-host only ... (nHosts=2). CI host microtests cover thenHostsgate, the empty-backendsupportedclear, and the NONE connection type.Not run here: the single-host SDMA+PROXY case. The new removal applies only when
nHosts > 1.Follow-ups
MixedBackends_TwoNodeUnsetTypewithNCCL_GIN_TYPEunset; today every 2-node config sets the type, so the test skips in CI.