Repository navigation
Conversation
On gfx1250 the per-peer Send/Recv decomposition that AllToAll falls back to is wrong in place (measured at every size >= 1 MiB, 8 and 16 ranks, eager and graph). In place now either takes a correct path or fails: - rcclSelectAlltoAll: in place keeps CE-Scratch wherever it fits ddaScratch even though the gfx1250 table does not force CE-Scratch (rcclAlltoAllCeScratchInPlace; RCCL_FORCE_CE=0 still turns it off). Otherwise ncclInvalidUsage with a WARN that names the API and the reason (rcclAlltoAllInPlaceGuard); a comm without a DDA fabric scratch (split / shrink / grow child, MNNVL off, multi-clique, DDA off) is named as such. The refusal goes through ncclGroupErrCheck, so it fails an enclosing group. Queries keep naming the fallback. - Registered windows (CTA_POLICY_ZERO): in place skips the registered-window CE AllToAll and hierarchical CE on gfx1250 (both race on the user buffer; rcclAlltoAllRegisteredCeRefused). - taskAppend: ncclAlltoAllConfig (no decision) gets the same CE-Scratch exception and the same refusal of the symmetric-window CE arm; the guard runs before the Send/Recv decomposition for every AllToAll. - NCCL_ENQUEUE_REARCH_ENABLE: guard before the lowering, which always decomposes AllToAll into Send/Recv. - CE-Scratch AllToAll on fabric comms: zero the staged scratch range and sync once more after the copy-back. The LL / LL128 kernels poll flag words in the same scratch with no entry barrier. - CHANGELOG entry. Other arches: unchanged (the helpers are gfx1250-only).
✅ 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.
🟡 Changes recommended
Check mode can incorrectly reject eligible CE-Scratch calls, and the LL128 alignment model diverges from production eligibility.
2 open findings
What changed in this PR
Guards gfx1250 in-place AllToAll against unsafe Send/Recv fallback, using CE-Scratch where possible and rejecting unsupported cases.
Changes:
- Adds shared in-place selection guards and CE-Scratch fallback.
- Clears and synchronizes fabric scratch after CE AllToAll.
- Adds host tests and release documentation.
| File | Description |
|---|---|
projects/rccl/src/rccl_wrap.cc |
Adds selection guards and warnings. |
projects/rccl/src/include/rccl_common.h |
Declares shared guard APIs. |
projects/rccl/src/enqueue/enqueue.cc |
Guards enqueue and rearchitected paths. |
projects/rccl/src/ce_coll.cc |
Clears scratch and adds synchronization. |
projects/rccl/test/host/wrap-test.cc |
Tests selection and warning behavior. |
projects/rccl/test/host/enqueue-test.cc |
Tests enqueue fallback handling. |
projects/rccl/test/host/ce-coll-test.cc |
Tests scratch cleanup ordering. |
projects/rccl/test/host/fakes/enqueue_test_deps.h |
Declares enqueue test seams. |
projects/rccl/test/host/fakes/enqueue_test_deps.cc |
Implements enqueue test fakes. |
projects/rccl/test/host/fakes/collectives_fakes.cc |
Adds enqueue-path test behavior. |
projects/rccl/CHANGELOG.md |
Documents the behavior change. |
🧠 Review effort: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
- taskAppend: an in-place AllToAll with a DDA fabric decision that NCCL_CHECK_MODE diverted to enqueue now takes CE-Scratch where it fits, like an undecided call, instead of being refused. The enqueue-site WARN now names the CE-Scratch reason and no longer cites suspended comms, which ncclEnqueueCheck rejects before taskAppend. - wrap-test: the LL128 eligibility fake no longer requires 16-byte aligned buffers (the V1 predicate ignores the pointers); cover a misaligned in-place call inside the LL128 window. - wrap-test: drive the selector refusal and the grouped CE-Scratch arm through ncclGroupStartInternal / ncclGroupEndInternal.
Conflict resolution: - rcclSelectAlltoAll: develop gates forced CE-Scratch on the ceNonReg size window (unless CTA_POLICY_ZERO); the in-place exception bypasses it, since CE-Scratch is the only correct in-place path there. - taskAppend: keep develop's merged AllReduce / ReduceScatter decided branch and its ReduceScatter exclusion from CE-Scratch, with the in-place AllToAll predicates on top. - Test fakes: develop moved the enqueue binary's stubs into per-TU fakes. The in-place AllToAll seams move to rccl_wrap_fakes; the taskAppend CE entry points keep failing loud unless a test sets g_rcclTaskAppendCeFaked. - CHANGELOG: the bullet moves to the new Unreleased section.
Conflict resolution: - rcclSelectAlltoAll: keep the gfx1250 in-place refusal of registered and hierarchical CE, with develop's rcclSelectDone() returns (#13098). - rccl_wrap_fakes: #13098 added functional rcclCeAllReduceEnabled / rcclForceCeAllReduceEnabled fakes, which now replace the taskAppend-gated ones (AllReduce-only calls, never reached by AllToAll), and calls rcclCeArNonRegCap from taskAppend, which had no fake; add a fail-loud one so rccl-UnitTestsMicroEnqueue links.
CHANGELOG conflict only: keep develop's Unreleased section and add this branch's gfx1250 in-place AllToAll bullet under Changed.
rccl_wrap_fakes: develop (#14633) removed the same duplicate CE AllReduce fail-loud fakes; keep this branch's taskAppend CE seams and the rcclCeArNonRegCap fake.
Develop (#14635) now fakes rcclCeArNonRegCap; drop this branch's copy.
i-kosarev
left a comment
i-kosarev
left a comment
There was a problem hiding this comment.
Corruption starts at ddaScratchBytes (256 MB) on a pre-#13336 build and at 1 MB after it, cause #13336 stopped forcing CE-Scratch. Host tests use fakes. This is a device timing race. Can we add into CI on-GPU in-place check with rank skew? It'd have caught issues uncovered by #13336. Could nccl.h.in document in-place ncclAlltoAll? The deprecated ncclAllToAll says unsupported. This PR now defines the semantics
There was a problem hiding this comment.
🟡 Changes recommended
A critical capture-detection gap lets DEBUG_LOCAL in-place AllToAll calls enter a graph-unsafe CE-Scratch path, risking replay hangs and requiring a fix with regression coverage.
1 open issue
2 resolved since last review
Review effort: Balanced
…ions - taskAppend: rcclSelectAlltoAll probes capture only on the way to its CE and Direct branches, so a DDA / GIN decision that NCCL_CHECK_MODE diverts here carries ceCapturing = false. Probe the stream for it, so a captured in-place call cannot take the graph-unsafe CE-Scratch path; it is refused instead. - rcclAlltoAllCeScratchInPlace: require nRanks > 1, like rcclAlltoAllRegisteredCeRefused. - nccl.h.in: document in-place ncclAlltoAll.
|
@i-kosarev thanks for the review.
|
Picks up #14579, which fixes the two one-RPN AllGather host tests.
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
There was a problem hiding this comment.
ARGUS
Findings
4 more finding(s), not posted inline
projects/rccl/src/rccl_wrap.cc:1139(Low) The selector-site WARN blames NCCL_CHECK_MODE, but check mode does not turn DDA off in the selector; drop the branch.projects/rccl/src/rccl_wrap.cc:1112(Low) The gfx1250 in-place predicate is spelled out three times; build CeScratchInPlace on RegisteredCeRefused.projects/rccl/src/include/rccl_common.h:294(Nit) Comment says "these two" but documents three helpers.projects/rccl/src/enqueue/enqueue.cc:4817(Nit) a2aRegCeRefused here means what rccl_wrap.cc calls a2aRegCeSkipped; rename it to a2aRegCeSkipped.
This PR changes 1,046 lines across 12 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: Worth settling before approval: CeAlltoAllDdaScratchMicrotest is not in test_categories_micro.yaml so CTest never runs it, and fabric CE-Scratch AllGather still skips the scratch zero and barrier at ce_coll.cc:3899.
Comment @argus review to re-run, or @argus re-review after pushing.
pr-rocm-systems-13461-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
pr-rocm-systems-13461-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
- Guard WARN: drop the selector-site NCCL_CHECK_MODE branch (check mode does not turn the DDA fabric paths off in the selector), and say the fabric paths need a per-peer size that is a multiple of 16 bytes rather than aligned buffers (no fabric eligibility check reads the pointers). - Build CeScratchInPlace, RegisteredCeRefused and the guard on one gfx1250 in-place predicate. - taskAppend: a2aRegCeRefused -> a2aRegCeSkipped, matching the selector. - Trim review tags and bench numbers from comments. - test_categories_micro.yaml: list CeAlltoAllDdaScratchMicrotest so the CTest tiers run it.
|
Argus findings not posted inline, all addressed in 51719d7:
Host microtests: Micro 3020/3020, MicroEnqueue 264/264, MicroCe 72/72, MicroInit 739/739, MicroTaskPrep 421/421. |
There was a problem hiding this comment.
ARGUS
This PR changes 1,029 lines across 12 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: Nothing new worth raising; still open from earlier: ce_coll.cc:3899. Worth clearing before approval.
Comment @argus review to re-run, or @argus re-review after pushing.
pr-rocm-systems-13461-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
pr-rocm-systems-13461-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
PR Pundit review(Testing PR Pundit Code Review agent. Not all comments may be relevant. This is a one-time automated pass: please address what is useful, no need to request a re-review. A quick reply per point (useful or not, and why) helps us calibrate the agent. Reach out to @peymanr for suggestions and feedback.)
1. 🟠 Should fix: In-place gfx1250 AllToAll hard-fails on split/shrink/grow comms and all sizes: this PR adds a refusal guard
On gfx1250, in-place Who is affected: gfx1250 applications that issue in-place AllToAll on split/shrink/grow children or other communicators without a fabric DDA scratch, or above the scratch size, or under capture. Checked against the code by an independent pass. Related past reviews: #9325 on Shared from a pilot of PR Pundit, an agentic code reviewer that builds and tests the change and reads the project's past reviews. Reviewed at |
|
@peymanr thanks. On PR Pundit #1 (in-place gfx1250 AllToAll refuses where no correct path exists): useful as a summary, but no change in this PR. It's the intended contract: on those communicators the only alternative is the per-peer Send/Recv decomposition, which has no ordering between "our send read byte b" and "our recv wrote byte b". So we return |
JIRA ID : AICOMRCCL-3440
Branch
users/thananon/gfx1250-a2a-inplace-guard, merged withdevelop4ed21d8. Standalone; independent of the gfx1250MNNVL registration / ncclMemAlloc PR and of the unroll-8 PR.
Summary
On gfx1250, in-place
ncclAlltoAllthat the DDA fabric LL / LL128 paths do not take falls back to per-peer Send/Recv, which is wrong in place (every size >= 1 MiB at 8 and 16 ranks, eager and graph; 773,991,356 wrong words on a 1 GiB misaligned in-place call).With this change an in-place AllToAll on gfx1250 either runs a correct path (DDA fabric LL / LL128, or CE-Scratch through
ddaScratch) or returnsncclInvalidUsagewith a WARN that names the API and the reason. It never decomposes into Send/Recv in place.The hazard is not gfx1250-specific. The in-place Send/Recv decomposition has no ordering between "our send has read byte b" and "our recv writes byte b" on any transport (P2P, net, graph-captured user buffers). Only MNNVL P2P at 8 / 16 ranks was measured wrong. "Other architectures are unchanged" means they keep the race; this PR fixes gfx1250 only, and the upstream issue should cover the p2p in-place decomposition, not only registered-window CE.
Change
rccl_wrap.cc,rccl_common.h: three gfx1250-only helpers shared by every AllToAll funnel.rcclAlltoAllCeScratchInPlacekeeps CE-Scratch for in place wherever it fits, although the gfx1250 table does not force CE-Scratch.rcclAlltoAllRegisteredCeRefusedkeeps in place off the registered-window and hierarchical CE AllToAll.rcclAlltoAllInPlaceGuardreturnsncclInvalidUsage+ WARN.rcclSelectAlltoAll: the CE-Scratch exception (in place it also ignores theceNonRegMin/ceNonRegMaxsize window that gates forced CE-Scratch), the registered-CE skip, and the guard before the p2p fallback (live calls only; queries keep naming the fallback). The refusal goes throughncclGroupErrCheck, so insidencclGroupStart/Endit fails the group, as thetaskAppend-site refusal does.ddaFabricMemHandlerit says "this communicator has no DDA fabric scratch (child comm from split/shrink/grow, MNNVL off, multi-clique, or DDA disabled)" instead of pointing at alignment.enqueue.cctaskAppend:ncclAlltoAllConfig(no decision) gets the same CE-Scratch exception and registered-CE refusal. The guard runs before the Send/Recv decomposition.rawTaskAppend(NCCL_ENQUEUE_REARCH_ENABLE) runs it before lowering.ce_coll.cc: on fabric comms, a CE-Scratch AllToAll zeroes the staged scratch range and syncs once more after the copy-back. The LL / LL128 kernels poll flags in the same scratch with no entry barrier, so a peer's next LL kernel could otherwise overwrite the scratch mid copy-back or accept stale payload as a flag. This is a separate, pre-existing hazard (CE-Scratch followed by an LL kernel, reachable on develop withRCCL_FORCE_CE=1/CTA_POLICY_ZERO). It is fixed here because this PR makes CE-Scratch default-reachable for in place. Eager only: CE-Scratch is never selected under graph capture.CHANGELOG.md: one "Changed" bullet.What changes on develop (gfx1250 in place only)
ddaScratch(8r ~546 MiB, 16r 1 GiB)RCCL_FORCE_CE=0, or under graph capturencclInvalidUsageNCCL_CHECK_MODEncclInvalidUsage(1 GiB)ncclAlltoAllConfigncclInvalidUsageNCCL_ENQUEUE_REARCH_ENABLEncclInvalidUsage(REARCH has no CE AllToAll)CTA_POLICY_ZERO+ registered windowsncclInvalidUsageNCCL_ALGOis ignored for the in-place CE-Scratch route (correctness over a hint), unlike the forced CE-Scratch arm, whichNCCL_ALGOstill blocks.By communicator type (gfx1250, in place, count > 0, nRanks > 1)
ddaScratchand the fabric LL / LL128 paths exist only on a root comm that passesncclDdaFabricCommInit(!parent && !isGrow,comm->MNNVL == 1, 2 <= nRanks <= 144, one clique, DDA scratch size > 0, VMM allocation succeeds). gfx1250 has no IPC DDA fallback.symmetricSupportRCCL_FORCE_CE=0, registered windows:ncclInvalidUsagelaunchOrderImplicit=1ncclInvalidUsage(the 64 KiB LL case included); a refusal failsncclGroupEndncclInvalidUsageNCCL_WIN_ENABLE=0, nosymmetricSupport)ncclInvalidUsagencclCommSplit/ncclCommShrink/ncclCommGrow) of AncclInvalidUsageat every size: no fabric handler, no scratchNCCL_MNNVL_ENABLE=0,NCCL_P2P_LEVEL=0,NCCL_CUMEM_ENABLE=0, build withoutHIP_FABRIC_API, fabric not ACTIVE/READY, zero UUID)ncclInvalidUsageat every sizencclInvalidUsageat every sizeRCCL_DDA_ENABLE=0,RCCL_DDA_FABRIC_BUFFER_SIZE=0, scratch OOM, nRanks > 144ncclInvalidUsageat every sizeRow B is the one most users would hit: PyTorch device-mesh subgroups are created with
ncclCommSplit. Whether torch on ROCm reachesncclAllToAllwithsend == recvis unverified. A staged fallback (copy the send region to a comm-owned temp, then out-of-place Send/Recv) would close B-E and is the planned follow-up.Large in-place and in-place under capture on row A stay
ncclInvalidUsageuntil a correct fast path covers them; nothing in this PR closes B-E.Out of place: unchanged, except the CE-Scratch AllToAll on fabric comms (explicit
RCCL_FORCE_CE=1/CTA_POLICY_ZERO) gains one T-byte memset and one CE barrier.Notes for reviewers
sendbuff == recvbuff. Overlapping but unequal buffers are undefined in NCCL and not detected.ncclAlltoAllvin place is out of scope: it is not guarded and has the same hazard.ENABLE_ROCSHMEM[_GIN], Pivot needs the Rome topology). None of them has an in-place answer, so the guarantee is stated for builds without them.RCCL_FORCE_CE=1/CTA_POLICY_ZERO; that is deferred.-Osays. A gfx1250 in-place shape with no correct path therefore failsalltoall_perfat warm-up; use-w 0 -O 1for such cells.Test plan
HIP_VISIBLE_DEVICES=-1: after merging develop (incl. test(rccl): assert RCCL's real one-RPN AllGather behaviour #14579): Micro 3020/3020, MicroEnqueue 264/264, MicroCe 72/72, MicroInit 739/739, MicroTaskPrep 421/421. New: in-place helpers, guard WARN per site / reason (incl. the no-fabric-scratch comm), selector refusal inside a group setsncclGroupErrorand failsncclGroupEndInternalwith nothing queued, grouped in-place CE-Scratch, selector CE-Scratch / registered-window / no-path cases, 12taskAppendAllToAll cases (incl. a captured, check-mode-diverted DDA decision that must not take CE-Scratch) (Config, REARCH, forced arm, NCCL_ALGO, decided fabric incl. the check-mode CE-Scratch route, registered windows), CE-Scratch trace (copy-back, memset, barrier on one stream). Mutation-checked.wrong= mismatching 32-bit words):ncclInvalidUsagencclInvalidUsagencclInvalidUsage(out of place captures fine, 0 wrong)NCCL_MNNVL_ENABLE=0, 4 MiB / 64 MiBncclInvalidUsagencclAlltoAllConfig32 MiB misaligned / 512 MiBncclAlltoAllConfig576 MiB (above the scratch)ncclInvalidUsageNCCL_CHECK_MODE32 KiB / 8 MiB / 64 MiB / 1 GiBncclInvalidUsageNCCL_ENQUEUE_REARCH_ENABLE=164 KiB / 32 MiBncclInvalidUsagencclAlltoAllConfigin place on the full gfx1250 AllToAll stack (same guard code): before this change 2,507,967 / 32,232,448 / 597,874,624 / 675,664,352 wrong at 32 MiB / 64 MiB / 512 MiB / 576 MiB; after, 0 wrong, and 576 MiB+ returnsncclInvalidUsage.Still owed (follow-up runs, not blocking review):
ncclGroupEnd->ncclInvalidUsage).CTA_POLICY_ZERO+ window-registered in place: needs the MNNVL registration PR to register windows on gfx1250, so it is validated there.ncclInvalidUsage+ the no-fabric-scratch WARN.Refs: an upstream issue for the in-place Send/Recv decomposition and registered-window CE AllToAll in place will follow.