Repository navigation
gin: keep the RMA proxy available when NICs are fused - #11133
Conversation
✅ 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 |
There was a problem hiding this comment.
🟢 Approval recommended
The change is narrowly scoped, aligns with the stated derivation semantics, adds clear diagnostic logging, and introduces targeted regression tests for the previously uncovered fused-NIC configuration.
Pull request overview
This PR fixes a capability-gating bug where NIC fusion incorrectly disabled the host RMA proxy for GIN communicators, adds a consistent init-time log line to expose the derived GIN/RMA support state, and introduces MPI tests to prevent regressions in fused-vNIC configurations.
Changes:
- Removes the
!globalNicFusedterm fromcomm->globalRmaProxySupportderivation so host RMA remains available under NIC fusion. - Adds an always-on init log line reporting the inputs/outputs of the GIN + host-RMA capability derivations.
- Adds two communicator-level MPI tests plus a dedicated test-runner config to validate
hostRmaSupportandginTypebehavior when vNIC fusion is present.
File summaries
| File | Description |
|---|---|
| projects/rccl/src/init.cc | Drops NIC-fusion from the host-RMA proxy gate and logs the full GIN/RMA support derivation inputs/outputs. |
| projects/rccl/test/transport/GinNicFusionMPITests.cpp | Adds MPI tests that require multi-node + fused vNIC and assert hostRmaSupport and ginType remain valid. |
| projects/rccl/test/CMakeLists.txt | Wires the new GinNicFusion MPI tests into the MPI unit test target. |
| projects/rccl/tools/scripts/test_runner/configs/net_ib_transport.json | Adds a net_ib_gin_nic_fusion runner config to execute the new tests with GIN enabled/type set. |
Review details
- Files reviewed: 4/4 changed files
- Comments generated: 0
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
There was a problem hiding this comment.
🟡 Changes recommended
The regression configuration is not scheduled, and rank-local test failures can deadlock peers.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Review details
Suppressed comments (1)
projects/rccl/test/transport/GinNicFusionMPITests.cpp:97
- If this local topology query fails on only one rank, that rank returns while every peer proceeds into
AnyRank(localFused), deadlocking the test until its 120-second timeout. Coordinate the result before any rank returns.
if (ncclTopoCheckNicFused(comm, &localFused) != ncclSuccess) {
ADD_FAILURE() << "ncclTopoCheckNicFused failed";
return false;
}
- Files reviewed: 4/4 changed files
- Comments generated: 3
- Review effort level: Balanced
52d1e65 to
bbcecd9
Compare
There was a problem hiding this comment.
🟡 Changes recommended
The fused transfer test is not selected by the runner and can deadlock when its rank-local operation fails.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Review details
- Files reviewed: 6/6 changed files
- Comments generated: 2
- Review effort level: Balanced
965575b to
5aadd7b
Compare
e5e5415 to
dc0858d
Compare
abouteiller
left a comment
abouteiller
left a comment
There was a problem hiding this comment.
Ready to merge after the minor fix to clarify documentation and changelog for effect on GIN
0ddb384 to
7d4f38d
Compare
|
%RCCL_TEST_CONFIG |
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
|
%RCCL_TEST_CONFIG |
|
%RCCL_TEST_CONFIG |
|
%RCCL_TEST_CONFIG |
7d4f38d to
c318c5d
Compare
c318c5d to
44cebe8
Compare
globalRmaProxySupport carried a !globalNicFused term, so hostRmaSupport went false for every multi-node communicator whose ranks had fused NICs. That took away the one-sided RMA API, the inter-node MR for symmetric memory, and the window returned by ncclCommWindowRegister. The term connects two things that do not touch. NIC fusion merges IB devices for the regular transport; GIN and RMA are registered with makeVDevice = NULL and select from their own topology node types, so they never receive a fused device. Upstream v2.30.7-1 resolves it the same way.
The derivation that lost the RMA proxy printed only a subset of its inputs, and
only when symmetric memory was already unavailable, so a run that lost the proxy
left no trace naming the cause. Emit all of them unconditionally:
GIN/RMA support: ginConnection 1 hostRma 1 (ginTypeMask 0x4 mloPart 0
nicFused 1 crossNic 1 railable 1 rmaPlugin 1 cuMemGdr 1 oneLsaTeam 0
numRmaCtx 1)
GinNicFusionMPITest brings up a multi-node communicator, confirms through ncclTopoCheckNicFused that a rank really fused a vNIC, and asserts on ncclCommQueryProperties. HostRmaSurvivesNicFusion is the discriminating test and fails on an unfixed build; GinTypeSurvivesNicFusion is an anti-regression guard, since ginType never depended on fusion. RmaMPIFusedNicTest covers the proxy's three send paths over a fused vNIC. The proxy is never handed a fused device in production, so this establishes that it degrades to one NIC of the pair rather than misbehaving. The fixture builds the vNIC once per process: makeVDevice appends to the shared device table and there is no destroyVDevice, so one per test would leave duplicates behind. RegMr allgathers base addresses and rkeys, so allocation is agreed across ranks before any registration and no registration is short-circuited by another.
net_ib_gin_nic_fusion runs the five unit tests above against the GIN proxy backend. gin-d3-nicfusion repeats that lane's gin-d3 device-kernel scenario with fusion on; the device-GIN path is identical either way, so it is a no-interference check rather than new coverage. Both name the NICs they merge instead of setting NCCL_NET_MERGE_LEVEL: the runner collects exit codes only, so a merge level that happened not to fuse would leave the suite green while testing nothing.
Both guides instructed NCCL_IB_MERGE_NICS=0, which contradicts the default and the behaviour this branch restores. GIN and the host RMA proxy are registered without a makeVDevice callback and select from their own topology node types, so they keep using physical devices and never receive a fused vNIC; fusion aggregates bandwidth for the regular transport only. Say that, rather than leaving users to disable the configuration the fix makes usable.
|
Override: RCCL PR gate on rocm-systems#11133 at 183302f. As per input CoCo (Allen) classifies the failure as perf collectives failing to run on ROCm 10.2. Run 38476550318 hits HIP failure 'invalid argument' in cuMem teardown after the collective has printed. The same signature on Ruby run 38384594630 was one single-node test, device-api-symmetric_reduce_perf, which does not exercise this PR’s fused-NIC change. Tracked: TheRock#9009. The earlier label failure is tracked separately by TheRock#8995 and rocm-systems#13268. |
TheRock Submodule Bump ActivityNewest first
|
Motivation
comm->globalRmaProxySupportcarried a!globalNicFusedterm that upstream dropped, sohostRmaSupportwent false for every multi-node communicator whose ranks had fused NICs. Such a rank losesncclPutSignal/ncclSignal/ncclWaitSignal(they returnncclInvalidArgument), the inter-node MR for symmetric memory, and the window fromncclCommWindowRegister, which returns success with*winNULL.The term connects two things that do not touch. NIC fusion merges IB devices for the regular transport; GIN and RMA are registered with
netInfo.makeVDevice = NULLand select from their own topology node types, so they never receive a fused device. For GIN, "usable when NIC fusion is enabled" can only mean coexistence, not exploitation — and it was the coexistence that broke, for the host half only.Changes
init.cc— drop the!globalNicFusedterm, matching upstream v2.30.7-1. CHANGELOG entry included.Plus one unconditional
INFOnaming every input to both derivations, replacing a conditional line that printed a subset only once something was already missing; and the docs that told users to setNCCL_IB_MERGE_NICS=0, contradicting both the default and this change.Tests:
HostRmaSurvivesNicFusionis the discriminating one and fails on an unfixed build;GinTypeSurvivesNicFusionis an anti-regression guard, sinceginTypenever depended on fusion;RmaMPIFusedNicTest×3 establish the proxy degrades to one NIC of the pair rather than misbehaving when handed a fused device, which production never does.Three runner lanes, all pinning
NCCL_NET_MERGE_LEVEL=PHBso they cannot skip green where the default level fuses nothing:net_ib_gin_nic_fusionnet_ib_rma_fused_gin_offnet_ib_host_api_fusedHostApiTestunder fusion, minusGraphCaptureReplayPutWaitThe fix is backend-independent: the derivation has no
ginTypeterm.Issue Tracking
JIRA ID : AICOMRCCL-1970
Test Result
Two nodes, eight RoCE NICs fused into vNICs. Every run carries 12
Made virtual device ... ndevs=2lines andnicFused 1in the capability line, so fusion is real rather than merely requested.HostRmaSurvivesNicFusionGinTypeSurvivesNicFusionhostRma!globalNicFusedrestoredBoth SKIP with
NCCL_IB_MERGE_NICS=0and withNCCL_GIN_ENABLE=0; re-measured on the current head. Fusion lane: 5 PASS. Host RMA API under fusion: 25 passed / 0 failed, identical without fusion.GinMPIDeviceTests: 41 passed / 0 failed with fusion, identical without; a 4-rank run adds the five barrier cases, 46 of 53.HostApiTest.WindowRegisterDeregisterpasses in all three of fusion+GIN-off, no-fusion+GIN-off and fusion+GIN-on.Not covered
HostApiTest.GraphCaptureReplayPutWaitfails multi-node on this stack —ncclPutSignalreturns an error under graph capture — identically with fusion, without fusion, and withNCCL_GIN_ENABLE=0, so it is independent of this change and warrants its own ticket.GIN backends other than the proxy were not run under fusion, and fusion cannot reach them:
ginRocshmemGdaDevicesandginAnvilDevicesboth hardcode*ndev = 1and enumerate no IB devices, andNCCL_GIN_TYPE=3is not built on ROCm.iflushon a fused device stays uncovered: everyIFlushAfterIGetparameterization fails inside the adapter's flush QP with and without fusion, and the case has no CI wiring.Submission Checklist