Repository navigation
fix(rccl): gate and bound the destroy-path intranode barrier - #12446
prasanna-amd wants to merge 13 commits into
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.
Copilot review overview
🟡 Changes recommended
The barrier can still deadlock during teardown, and the new peer lookup can access memory outside the array.
Review effort: Balanced
Findings: 2
Open (4)
What changed in this PR
This PR changes RCCL communicator teardown to avoid a bootstrap barrier when a local rank was excluded by shrink or when the communicator is single-node.
Changes:
- Track whether shrink excludes a host-local rank.
- Use that flag and the node count to decide whether teardown enters the barrier.
- Add host tests for the new flag and update shrink-test peer data.
| File | Description |
|---|---|
| projects/rccl/test/host/init-test.cc | Updates shrink fixtures and adds host-local exclusion tests. |
| projects/rccl/src/init.cc | Sets the exclusion flag and changes the teardown barrier guard. |
| projects/rccl/src/include/comm.h | Stores the new flag on the communicator. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: nNodes > 1 also removes the only ordering between the legacy IPC deregister drain at init.cc:4026-4032 and ncclProxyStop, and that ordering is host-local, so nNodes == 1 is its worst case.
Comment @argus review to re-run, or @argus re-review after pushing.
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: comm->peerInfo can still be null at init.cc:4540, since ncclCommEnsureReady returns success out of its abortFlag branch without checking initState.
Comment @argus review to re-run, or @argus re-review after pushing.
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: comm->nNodes > 1 is not the same test as "no host-local proxy dependency", so a single-node comm relaying through a peer's proxy skips the barrier at init.cc:4033.
Comment @argus review to re-run, or @argus re-review after pushing.
fc0050a to
f2dfec7
Compare
The NCCL v2.31.2-1 sync added bootstrapIntraNodeBarrier in commDestroySync to protect PXN proxy teardown. This barrier requires all host-local ranks to participate, which deadlocks in two cases: 1. After ncclCommShrink: excluded ranks abort and skip the barrier while non-excluded ranks enter it (AICOMRCCL-2467). 2. Single-node configs: PXN is not active (nNodes == 1) but processes reach destroy independently and the barrier deadlocks (AICOMRCCL-2468). Guard the barrier on childCount == 0 (no shrink/split children) and nNodes > 1 (PXN is multi-node only).
Replace the childCount == 0 condition with a dedicated hasExcludedLocalRank flag. childCount is never incremented for shrink, so the old guard did not skip the barrier for shrink parents. It also incorrectly removed PXN teardown protection after normal split/grow operations. The new flag is set during shrink when any excluded rank shares the same host (via hostHash comparison). This precisely targets the case where a local peer will skip the barrier after abort, while preserving the barrier for healthy multi-node comms including those with hierarchical splits.
…ield comment Move peerInfo_ array into Grow_ParentComm so all shrink tests have valid peerInfo without per-test setup. This fixes the null dereference in ShrinkIgnoresNoColor and eliminates copy-pasted peerInfo boilerplate across five tests. Add inline comment to hasExcludedLocalRank field.
…de test Add null check on comm->peerInfo before the hostHash loop — a non-blocking comm whose init failed before allgather can have peerInfo null with abortFlag set. Give Grow_ParentComm entries distinct non-zero hostHashes so existing shrink tests don't accidentally match on zero-initialized hashes. Add a multi-element exclude test where the first entry misses and the second hits, exercising loop continuation.
Setting the flag for any shrink was not merely a lost optimization, it introduced a new hang. Under NCCL_SHRINK_DEFAULT the excluded rank is healthy: its abortFlag is 0 and it still destroys its parent, so it does enter bootstrapIntraNodeBarrier. A surviving host-local rank that skipped the barrier would strand it there forever. Only NCCL_SHRINK_ABORT leaves behind a peer that will never arrive. The three host tests covering this path were written against NCCL_SHRINK_DEFAULT, which was itself the tell; they now use NCCL_SHRINK_ABORT. A fourth asserts NCCL_SHRINK_DEFAULT leaves the flag clear so the regression cannot return quietly. Also adds permanent coverage for the destroy-path barrier deadlock to TeardownStressTests.cpp. Two ranks, one process each, divergent teardown: the control has both ranks destroy, the regression has the peer abort so it skips the barrier. Both run under ProcessIsolatedTestRunner with a hard timeout, so a regression fails the test rather than wedging CI.
The barrier exists only to protect PXN connection establishment, so test that directly instead of using nNodes > 1 as a proxy for it. nNodes > 1 was not the same question. proxyRank at net.cc:489 comes from the graph == NULL branch of ncclTopoGetNetDev(), which under the default NCCL_P2P_PXN_LEVEL=2 assigns a peer's rank from purely local topology -- no nNodes test, no system->inter test. The earlier justification for nNodes > 1 cited ncclTopoGetPxnRanks(), but that function only feeds the proxy pre-connect at init.cc:2511 and does not gate net.cc at all. A single-node comm with NCCL_PXN_DISABLE=0 really can relay through a host-local peer's proxy. ncclPxnDisable(comm) answers the actual question and every rank agrees on it, since it derives from arch, nRanks and env. Agreement is the property that matters: a barrier some ranks enter and others skip is what hangs. comm->useNetPXN would be more precise still, but it is unusable here -- the allgather that makes it comm-wide is conditional on LOCAL_REGISTER or GRAPH_REGISTER (generic.cc:30), and net.cc:489 keeps writing it afterwards for lazily established connections, so ranks can disagree at destroy. Also harden ncclPxnDisable() against a null comm->topo, now that it is reachable from the teardown path, and rename hasExcludedLocalRank to hasExcludedHostRank: it is computed from peerInfo[].hostHash, while every other local* field in ncclComm is node-scoped, and node != host on MNNVL.
EnvVars reports 0 GPUs inside a process re-exec'd by ProcessIsolatedTestRunner and a skipped child is scored as a pass, so the maxGpus gate was silently turning both divergent-teardown tests into no-ops that reported PASS without ever forking a rank. Guard with !isIsolatedChild(), matching TeamTopologyTests, and read the count from a cached EnvVars() instead of constructing a TestBed before fork(). Also drop the retracted ncclTopoGetPxnRanks justification from the header comment; the reason a two-rank single-node comm has no PXN relay is RCCL's NCCL_PXN_DISABLE=1 default plus rcclSetPxn()'s rank thresholds.
327297f to
e66003d
Compare
Review feedback on #12446: skipping the host-local barrier in commDestroySync() whenever a shrink had excluded a same-host peer threw away PXN teardown protection for every surviving rank, when only the dead peer actually needed to be left out. Record the ranks an NCCL_SHRINK_ABORT removed, and drop just those from the barrier membership so the survivors still rendezvous. Every survivor holds the same exclude list, so they derive the same membership and the same root, and bootstrapP2PBarrier already returns early at nranks == 1. Also fix the PXN gate polarity: ncclPxnDisable() returns an env-derived value the rest of the codebase compares against 1, so == 0 wrongly skipped the barrier for NCCL_PXN_DISABLE=2.
|
Pushed Barrier membership instead of skipping it. PXN gate polarity. Comment layout. Upstream's comment is restored verbatim inside the Field placement. The new fields live in the Tests. Host tests reworked to assert the recorded set rather than a bool, with a non-zero caller rank, accumulation across two shrinks, and a direct test of the membership helper. 707/707 Still out of scope and tracked separately: with PXN genuinely enabled, an |
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: a rank excluded by NCCL_SHRINK_ABORT that destroys its parent rather than aborting now enters the barrier the survivors filtered it out of, and hangs where it previously completed.
Comment @argus review to re-run, or @argus re-review after pushing.
| } | ||
| int const peer = comm->localRankToRank[i]; | ||
| if (comm->peerInfo[peer].hostHash != comm->peerInfo[comm->rank].hostHash) continue; | ||
| if (peer != comm->rank && commRankWasAbortExcluded(comm, peer)) continue; |
There was a problem hiding this comment.
Argus already flagged this and it is still open. An excluded rank never ran shrink, so destroying its own parent or a sibling comm sends it back into the barrier waiting for everyone, while the survivors already left. This filter does not cover that, so multi-node still hangs. How do we close it, and would a barrier with a timeout be the fuller fix?
There was a problem hiding this comment.
Yes -- a barrier with a timeout is the fuller fix, and that is what b7bc5b6 does. The filter is gone; hostRanks is upstream's loop verbatim again.
Your framing is what made the difference. Filtering tries to predict who will not arrive, which is only possible for the one failure mode we happened to be looking at, and only from the side that happens to know. A deadline does not need to know who or why. It covers the excluded rank destroying its parent, it covers a plain ncclCommAbort on a host-local peer (the AICOMRCCL-2468 shape I had previously said needed a follow-up notification protocol), and it covers a rank that simply died.
Mechanically:
ncclSocketAccepttakesint64_t deadlineNs = 0. It is a busy spin oversocketProgressState, not a blockingaccept(2), so the deadline drops straight into the existing loop. The expiry check must return before the post-loopswitch, which maps a still-Acceptingsocket toncclSuccessfor theasyncFlagcallers -- otherwise a peer that never arrived reports as a completed accept.- Threaded through
bootstrapAccept->socketAccept->bootstrapRecvDeadline->bootstrapP2PBarrieras a defaulted parameter, so no existing caller changes behaviour. The connect side was already bounded (~59.5s retry ->ncclRemoteError), so the accept side was the only gap. - One deadline for the whole barrier, not one per round.
bootstrapIntraNodeBarrierTimeoutwraps it;RCCL_DESTROY_BARRIER_TIMEOUT_MSECtunes it and0restores the unbounded wait.
Expiry is non-fatal: commDestroySync already logged the barrier result and continued teardown, so a deadlock degrades into a late PXN connection failure on the peer that never showed. A rank that gives up also stops feeding its own dst, so its peers hit their own deadlines in turn and the host drains in roughly one timeout rather than cascading.
One thing I would like a second opinion on: the default is 120000 ms, which is my own pick and not derived from anything. It wants to be comfortably longer than a slow-but-healthy PXN connect and comfortably shorter than whatever job-level watchdog sits above us. If you have a better-grounded number, I will take it.
There was a problem hiding this comment.
Yes, and that is what the PR is now -- the filter is gone entirely in b7bc5b61c28, hostRanks is upstream's loop verbatim, and the deadline is the fix. Your framing is what moved it; filtering tries to predict who will not arrive, which only works for the one failure mode we happened to be looking at and only from the side that happens to know.
End-to-end coverage is now up (38d7a5b5e5d, DivergentTeardown_PeerAborts_PxnEnabled_TimesOut), so the mechanism is no longer asserted without a test.
Leaving this open for the one thing I still owe you an answer on: the 120000 ms default. I said it was my own pick and not derived from anything. I can now at least bound it from both sides, though I would still rather have your read than ship my own number unchallenged:
- Upper bound. PyTorch's
ProcessGroupNCCLwatchdog is 600s. The destroy deadline must fire well below that or it is never the thing that fires, and the user sees the watchdog's much less specific failure instead. - Lower bound. This wants to exceed any legitimate skew between host-local ranks of one comm reaching
ncclCommDestroy. That is sub-second in normal operation and seconds under a pathological stall, so 120s is orders of magnitude of headroom. - What a false positive costs. A peer merely slow rather than gone loses whatever PXN connection it was establishing. Non-fatal but real, which argues for generous over tight.
That reasoning admits anything from roughly 60s to 300s. I picked the middle. If you have a number grounded in an actual observed PXN connect time I will take it over mine.
| // Pinned to NCCL_PXN_DISABLE=1, RCCL's default. With PXN genuinely enabled an | ||
| // ncclCommAbort peer still strands the barrier -- that needs a cancellation | ||
| // protocol and is tracked separately; do not widen this case without one. | ||
| TEST(Teardown, DivergentTeardown_PeerAborts_DoesNotHang) |
There was a problem hiding this comment.
This test pins NCCL_PXN_DISABLE=1 and calls ncclCommAbort with no shrink, so it passes via the PXN skip, not via the filtering. The main new logic of this PR has no end-to-end test. Can we add a real shrink(SHRINK_ABORT) + destroy repro with PXN on that actually hits the filter?
There was a problem hiding this comment.
The filtering this was asking about is gone as of b7bc5b6, so the specific test you asked for -- shrink(SHRINK_ABORT) + destroy with PXN on, hitting the filter -- no longer has anything to hit.
But I do not want to close this out on that technicality, because the underlying complaint survives the revert: the main new logic of this PR still has no end-to-end test. It is now the timeout rather than the filter, and the gap is the same one you identified. What exists today:
DivergentTeardown_PeerAborts_DoesNotHangpinsNCCL_PXN_DISABLE=1, so it passes via the PXN skip and never reaches the barrier -- exactly your original point, still true.DivergentTeardown_PeerDestroys_PxnEnabled_CompletespinsNCCL_PXN_DISABLE=0and does run the barrier, but both ranks arrive, so it proves the barrier still works rather than proving the deadline fires.
What is missing is the case that matters: NCCL_PXN_DISABLE=0 plus one rank aborting, asserting that the survivor returns from ncclCommDestroy in about the timeout rather than hanging, and with a non-success rc so the assertion is falsifiable. That is the test that would have failed before this PR and passes after it. I am writing it next, along with host-level coverage of the timeout value actually reaching bootstrapIntraNodeBarrierTimeout and of an expiry not failing teardown.
One caveat I would rather state than have you discover: I cannot run GPU tests from my current checkout -- the compute nodes on this cluster have no path to it -- so the end-to-end case will be verified in PR CI, not locally. I will post the run link rather than just asserting it passes.
Holding this thread open until that is up.
There was a problem hiding this comment.
The test is up, in 38d7a5b5e5d: DivergentTeardown_PeerAborts_PxnEnabled_TimesOut.
NCCL_PXN_DISABLE=0 plus one rank calling ncclCommAbort. Rank 0 genuinely enters the barrier, the peer never arrives, and only the deadline can release it. Before this PR that hung forever.
Two details aimed at the vacuous-pass problem you raised, since "the test passes" is worth very little on its own here:
RCCL_DESTROY_BARRIER_TIMEOUT_MSECis pinned to 5000 in the child. At the 120000 default the runner's own timeout fires first and the case reports a hang rather than the deadline working.- Destroy must take at least 4000ms. This is the part that keeps it honest: a prompt return would mean the barrier was skipped via the PXN gate, so the deadline was never exercised -- precisely the "passes via the skip, not via the new logic" failure you called out originally. Without this floor the case would still go green if someone later widened the gate.
Expiry reports ncclRemoteError, which is correct rather than a failure, so this case asserts on duration instead of rc. The other two keep the strict rc == ncclSuccess check from pvallem's thread.
Also removed a comment above DivergentTeardown_PeerAborts_DoesNotHang that still claimed the PXN-on abort case "needs a cancellation protocol and is tracked separately". That was true of the old design and would have warned the next person off writing exactly this test.
Caveat I stated earlier and will restate: I cannot run GPU tests from my checkout, so this is verified by syntax check locally and must be proven in CI. I am not claiming it passes until I can post the run. CI has not run on this branch since 2026-10-08 -- the ci:single-arch label gate explains part of that but not all of it, and I am chasing the rest. Holding this thread open until the run link is here.
|
I think its scope has moved beyond AICOMRCCL-2467: the ticket concerns a caller-side c10d bug, and the PyTorch fix makes the test pass without any RCCL changes. The shrink-abort filtering (abortExcludedRanks) addresses a separate issue and should have its own ticket, I think. The new GPU test, DivergentTeardown_PeerAborts, does not exercise that filtering: it sets NCCL_PXN_DISABLE=1 and never calls shrink, so it passes because the PXN gate skips the barrier. Could you clarify what remains to be addressed and the goal of this PR, move the filtering work to a separate ticket. |
Revert the abort-excluded-rank filtering from 7118726 and replace it with a deadline on the barrier itself. Filtering was unsound. The exclude list exists only in the surviving processes, so a rank removed by NCCL_SHRINK_ABORT never learns it was dropped: it still holds its parent handle, still enters the full barrier, and waits for peers that already rendezvoused without it. That moves the hang onto the excluded rank rather than removing it, and regresses the GrowMPITests Grow_ThenShrinkAbort shape wherever PXN is on. Any membership derived from survivor-only knowledge has this flaw, including the earlier hasExcludedHostRank skip. What is left is two layers that both hold without agreement between ranks: 1. The PXN gate, unchanged from baf39df. The barrier protects PXN connection establishment only, and ncclPxnDisable() is derived from arch, nRanks and env, so every rank computes it identically. The hostRanks loop is restored to upstream verbatim. 2. A deadline. A peer can fail to arrive for reasons no membership calculation can predict - ncclCommAbort elsewhere in the process, a rank that died outright - so the barrier now gives up and reports an error rather than blocking forever. Expiry was already non-fatal here: the result is logged and teardown continues, degrading a deadlock into a late PXN connection failure on the peer that never showed. The deadline is plumbed as a defaulted parameter through ncclSocketAccept, bootstrapAccept, socketAccept and bootstrapRecvDeadline, so no existing caller changes behaviour. One deadline covers the whole barrier rather than each round. A rank that gives up stops feeding its own dst, so its peers hit their own deadlines in turn and the host drains within roughly one timeout instead of cascading. RCCL_DESTROY_BARRIER_TIMEOUT_MSEC tunes it; 0 restores upstream's unbounded wait. comm.h is back to no change at all against develop.
Three claims in the comments added by b7bc5b6 did not survive being checked. "One deadline for the whole barrier ... so the total wait is what the caller asked for" was false. bootstrapSend keeps its own connect retry per round, so the worst case is the deadline plus that retry. Says so now. "Degrades a deadlock into a late PXN connection failure on the peer" was asserted, not verified. It happens to be true, and now carries the evidence: the proxy service thread force-closes its peer and listen sockets on stop (proxy.cc:2090-2095, 2108), so the peer gets a closed socket or ECONNREFUSED rather than hanging. Worth recording because the opposite was plausible -- ncclProxyStop does call ncclShmClose on the proxy ops pool, which looks like it could unmap a segment peers still hold. It cannot: a client attach leaves handle->shmPath NULL and ncclOsShmClose unlinks only when it is non-NULL (os/linux.cc:777, 892), so only this rank's own view is released. Teardown is also refCount- gated (proxy.cc:2284), so sibling comms sharing the proxy are untouched. "A rank that died outright" was listed as a case needing the timeout. It is not: the kernel closes a dead process's sockets, so the recv already errors out. The case that actually needs this is a peer that is alive but not participating, having taken the abortFlag branch and skipped the barrier. That also explains why the ack read in socketAccept is left unbounded -- reaching it means a peer completed the TCP connect, which a dead process cannot do. The consequence worth naming: an alive-but-not-participating peer is indistinguishable from a slow one, so expiry is a guess, and a wrong guess is paid for by a healthy peer. That makes the default a correctness-relevant choice rather than a cosmetic one, so it is now justified in place -- 120s is far above the expected skew between host-local ranks reaching destroy, and below PyTorch's 600s watchdog, beyond which it could never be the thing that fires. No functional change.
prasanna-amd
changed the title
The deadlineNs parameter added to ncclSocketAccept left the host-test link stub at the old three-parameter form. Default arguments are not part of a function's type, so bootstrap.cc's calls resolved to a symbol nothing defined and rccl-HostUnitTests failed to link.
Five microtests over commDestroySync: the barrier runs with the configured deadline when PXN is on, is skipped when PXN is off or the comm is aborting, and on expiry still stops the proxy while reporting the error. Each was checked against a mutated guard to confirm it fails for the right reason. Also resets the recording fake between tests, and drops the peerInfo scaffolding in Grow_ParentComm left over from the reverted membership filter.
The PXN-on + peer-abort case is the one the gate does not cover: rank 0 genuinely enters the barrier and only the deadline can release it. Pins RCCL_DESTROY_BARRIER_TIMEOUT_MSEC low so the runner does not time out first, and requires destroy to take at least that long -- a prompt return would mean the barrier was skipped and the deadline never exercised. Also drops a comment that still claimed this case needed a cancellation protocol, which the deadline supersedes.
|
@nikitaxgusev -- overdue reply to your questions from 10-08. Taking them in order, because two of them dissolved when the design changed and two did not. "Move the filtering work to a separate ticket." No ticket needed -- the filtering is gone. Reverted in "DivergentTeardown_PeerAborts does not exercise the filtering." Correct then, and the more durable form of the complaint -- that the main new logic had no end-to-end test -- was still true after the revert, just about the deadline rather than the filter. Fixed in "Who will review the commDestroySync diff against next upstream during the next NCCL sync?" A fair worry, and the reason the hunk is now shaped the way it is. Upstream's three comment lines are back verbatim inside the "The ticket concerns a caller-side c10d bug, and the PyTorch fix makes the test pass without any RCCL changes." This is the one I cannot close out, and I would rather say so than hand-wave it. I have not verified it -- you did not cite a PR number and I have not found the change you mean. Could you point me at it? Two things I would want to separate before concluding it makes this PR unnecessary:
So I would frame the PR as: it does not depend on the 2468 attribution and should not be judged on it. If your PyTorch fix does resolve the c10d case, that is good and orthogonal -- happy to retitle to drop the 2467 framing if the ticket is genuinely caller-side. |
|
Withdrawing the question I just asked -- I found it: pytorch/pytorch#200372, "[c10d] Abort the parent group in shrink_group with SHRINK_ABORT" (open, fixes pytorch#194679). @nikitaxgusev, you were right and I should not have made you supply the reference. Its description names the mechanism independently of anything in this PR:
That is the same barrier and the same arrival asymmetry as the gdb trace in the description, reached from the caller side. So I accept your point: this PR should not be justified on AICOMRCCL-2467. The c10d-specific hang has a caller-side fix and does not need RCCL to change. What I think survives, and what I would like the PR judged on:
Concretely: happy to retitle and refile this against a new ticket for the general unbounded-barrier case, and let AICOMRCCL-2467 be closed as caller-side once #200372 lands. Your call on whether that is cleaner than keeping the existing number -- say which and I will do it. (AICOMRCCL-2468 is unaffected either way. Its attribution remains unproven and open; that run fails at |
…13789) JIRA ID: AICOMRCCL-2468 ## Summary The `whl-multi-arch` index on `rocm.nightlies.amd.com` **stopped publishing torch on 2026-08-22**. It still returns 200 and still serves wheels, so nothing failed — `uv` just resolves to the newest wheel the index offers, which has not moved since August. Every `pytorch_c10d` CI run since 2026-09-12 has installed the identical build: ``` 2.15.0a0+rocm10.1.0a20260822 ``` So for two months this job has been reporting green against a two-month-old PyTorch while appearing to test nightlies. This points the install at the current publishing location, `nightly.repo.amd.com/rocm/pytorch/whl-next/`, which today resolves `2.16.0a0+rocm10.2.0a20261019`. ## Why the extra index is required The `device-gfx942` extra resolves `rocm-sdk-device-gfx942`, which on the new host is published to `rocm/core/whl-next/`, **not** `rocm/pytorch/whl-next/`. With a single `--index-url` the resolve fails outright: ``` And because you require torch[device-gfx942], we can conclude that your requirements are unsatisfiable. ``` With both indexes it resolves cleanly (verified via `uv pip install --dry-run`): ``` + torch==2.16.0a0+rocm10.2.0a20261019 + rocm-sdk-core==10.2.0a20261019 + rocm-sdk-device-gfx942==10.2.0a20261019 + rocm-sdk-libraries==10.2.0a20261019 + triton==3.9.0+gitfc8d31a5.rocm10.2.0a20261019 ``` ## Test sources follow automatically `projects/rccl/ci/scripts/test_pytorch_c10d.py::clone_pytorch_test_sources` date-pins the cloned PyTorch test tree off the installed wheel's version string (the `--shallow-since` path). The test tree therefore moves forward with the wheel — no separate bump needed here. ## Relationship to AICOMRCCL-2467 `test_shrink_group_flags` has been timing out at 300 s in CI. The upstream teardown fix is pytorch/pytorch#194775 (landed 2026-10-13), which the frozen August wheel cannot possibly contain. Because the test sources are date-pinned off the wheel, this change pulls in both the fixed runtime and the updated test. This is complementary to #12446 (the RCCL-side destroy-path barrier fix), not a replacement for it. ## Deliberately floating, not pinned This tracks the latest nightly rather than pinning an exact build, which restores the originally intended behaviour of this job. The tradeoff is that an upstream nightly regression can turn CI red; that is preferred here over silently testing stale wheels, which is the failure mode this PR exists to fix. ## Test plan - [ ] `therock-rccl-test-pytorch-distributed.yml` dispatched on this branch, `test_scope: all` - [ ] Confirm the run log reports a `2.16.0a0+rocm10.2.0a2026101x` version rather than `...20260822` - [ ] Confirm `test_shrink_group_flags` no longer times out --------- Co-authored-by: renjithravindrankannath <94420380+renjithravindrankannath@users.noreply.github.com>
|
Closing: not needed. The test_shrink_group_flags hang (AICOMRCCL-2467) is fixed on the PyTorch side by pytorch/pytorch#194775 (aborts the shrink parent instead of destroying it), and the related teardown hangs by pytorch/pytorch#200673. No RCCL change is required for the reported failures. |
JIRA ID: AICOMRCCL-2467
Problem
The NCCL v2.31.2-1 sync added
bootstrapIntraNodeBarriertocommDestroySync()to protect PXN proxy teardown. It is a socket-based dissemination barrier that requires every host-local rank to participate, and it has no cancellation protocol. A rank that does not arrive strands every peer that did, forever.Ranks arriving at different times is fine — the barrier handles that correctly. Divergent participation is what hangs it, and the guard immediately above the barrier is itself a source of divergence: a rank with a non-zero
abortFlagskips the barrier entirely. SoncclCommAborton one rank — including the PyTorchProcessGroupNCCLwatchdog path, and anncclCommShrink(NCCL_SHRINK_ABORT)that excluded a host-local rank — leaves its peers blocked inbootstrapRecvwith no way out.In the reproduced configuration the barrier was protecting nothing. RCCL defaults
NCCL_PXN_DISABLE=1, andrcclSetPxn()only auto-enables PXN atnRanks >= 64(gfx942) />= 32(gfx950), so a single-node comm has no PXN relay unless PXN is forced on explicitly.Evidence
Captured under gdb from a live hang on a two-rank, single-node, one-process-per-rank reproducer:
ncclCommDestroyafter a 10s staggerncclCommAbortThe passing stagger control is what makes this meaningful: the barrier is sound when participation agrees.
On AICOMRCCL-2468: this proves the mechanism is real and reachable on a single-node comm. It does not establish that
DistributedDataParallelTest.test_nccl_backend_4gpu_modulefails this way — that run's last log line isInit START, which points at init rather than destroy. That attribution is tracked on AICOMRCCL-2468 and remains open.Fix
Two independent changes, both in
commDestroySync().1. Only run the barrier when PXN is actually enabled
The barrier exists solely to protect PXN connection establishment, so test that directly.
ncclPxnDisable()derives from arch,nRanksand env, so every rank agrees on it by construction — which is the property that matters. On the overwhelming majority of RCCL comms this removes a deadlock risk that was protecting nothing.comm->useNetPXNwould be more precise, but it is unusable as a guard: the allgather that makes it comm-wide is conditional onLOCAL_REGISTER/GRAPH_REGISTER(generic.cc:30), andnet.cc:489keeps writing it afterwards for lazily established connections. Ranks can therefore disagree at destroy, which would reintroduce the hang in a new shape.ncclPxnDisable()is now hardened against a nullcomm->topo, since the destroy path can reach it on a non-blocking comm whose init failed before the topology was built.2. When it does run, bound it
bootstrapIntraNodeBarrierTimeout()gives the barrier an absolute deadline (RCCL_DESTROY_BARRIER_TIMEOUT_MSEC, default 120000; 0 restores upstream's unbounded wait). On expiry it returnsncclRemoteError.commDestroySync()logs that, continues teardown --ncclProxyStop()still runs -- and propagates the error as its own result, which is upstream's existing handling for a failed barrier.No membership calculation can predict every way a peer fails to arrive — most of all an
ncclCommAbortelsewhere in the process, which makes that rank take theabortFlagbranch and skip the barrier. The deadline covers that class generically.The case this targets is specifically a peer that is alive but not participating. A peer that actually dies needs no timeout: the kernel closes its sockets on process exit, so the recv already errors. That asymmetry is why only the connect wait is bounded and the ack read in
socketAccept()is left alone — reaching it means a peer completed the TCP connect, which a dead process cannot do.What expiry costs is bounded. Proceeding means running
ncclProxyStop()while a peer might still be establishing a PXN connection through us — the race the barrier exists to prevent — so it is worth being precise. The proxy service thread force-closes its peer and listen sockets on stop (proxy.cc:2090-2095, 2108), so the peer gets a closed socket orECONNREFUSEDand fails cleanly. Nothing shared is unmapped out from under it: thencclShmClose()inncclProxyStop()releases only this rank's client-side view, because a client attach leaveshandle->shmPathNULL andncclOsShmClose()unlinks only when that is non-NULL (os/linux.cc:777, 892). Teardown is also refCount-gated (proxy.cc:2284), so comms sharing this proxy state are untouched.The residual risk is a false positive: a peer merely slow to reach destroy, rather than gone, loses whatever it was establishing. Hence a generous default rather than a tight one — 120s is orders of magnitude above the legitimate skew between host-local ranks of one comm reaching
ncclCommDestroy, and sits below PyTorch's 600sProcessGroupNCCLwatchdog so it is never the deadline that fires first.What this PR no longer does
Earlier revisions tried to vary the barrier's membership — a
hasExcludedHostRankflag, then anabortExcludedRanksfilter — to drop ranks that anNCCL_SHRINK_ABORThad excluded. Both were unsound and are fully reverted. Only the shrink survivors know about the exclusion; the excluded rank still holds its own parent handle and enters the full barrier, so filtering does not remove the hang, it relocates it onto the excluded rank.src/include/comm.his back to a zero-line diff againstdevelop.Thanks to @nikitaxgusev for pushing on this, and to amd-argus for catching that the earlier
nNodes > 1guard rested on a wrong reading ofncclTopoGetPxnRanks().Test
test/TeardownStressTests.cpp— four process-isolated GPU tests, each with a hard timeout so a regression fails rather than wedging CI:DivergentTeardown_PeerDestroys_Completes— control, both ranks destroy, PXN off.DivergentTeardown_PeerDestroys_PxnEnabled_Completes— same withNCCL_PXN_DISABLE=0, so the barrier genuinely runs.DivergentTeardown_PeerAborts_DoesNotHang— peer aborts and skips the barrier, PXN off, so this covers the gate.DivergentTeardown_PeerAborts_PxnEnabled_TimesOut— the case the gate does not cover:NCCL_PXN_DISABLE=0plus an aborting peer, so the barrier is entered and only the deadline releases it.RCCL_DESTROY_BARRIER_TIMEOUT_MSECis pinned to 5000 so the runner does not time out first, and destroy must take at least 4000 ms — a prompt return would mean the barrier was skipped and the deadline never exercised. Expiry reportsncclRemoteErrorhere, which is the expected outcome, so this case asserts on duration while the other three keep the strictrc == ncclSuccesscheck.NCCL_PXN_DISABLEis pinned per child viasetenvrather than inherited:TestConfiginherits the parent environment, and an ambientNCCL_PXN_DISABLE=0(mi300_ainic_ccs_roce.jsonsets it) would otherwise silently change which path is under test.test/host/init-test.cc— five microtests overcommDestroySync, against a recording barrier fake:ncclProxyStop()still runs.Each was checked against a mutated guard — dropping the PXN condition, and swapping the bounded call back for the unbounded one — to confirm it fails for the intended reason rather than passing vacuously.
Host unit suites: 708/708 (
rccl-UnitTestsMicroInit) and 317/317 (rccl-HostUnitTests) PASS.Still outstanding
7118726b2f8— the reverted design. Every commit making up the current fix has zero CI coverage. Theci:single-archlabel gate ontherock-ci.ymlexplains part of that; it does not explain why the ungatedrccl-host-unit-testsandrccl-default-behaviorworkflows also stopped.