Repository navigation
feat(rccl): Enable hierarchical CE scale-out collectives - #10450
Merged
thomas-huber merged 19 commits intoSep 18, 2026
Merged
Conversation
Route multi-LSA AllGather and AllToAll through the hierarchical RMA plus CE path, and scope synchronization to the local LSA. Correct proxy window handles, offsets, and scratch cleanup so data-bearing RMA operations complete safely. AICOMRCCL-630 Co-authored-by: Cursor <cursoragent@cursor.com>
Add topology-gated multi-node AllGather and AllToAll coverage with repeated data-distinct iterations and hierarchical path assertions. Keep white-box synchronization expectations aligned with the LSA-local rank domain. AICOMRCCL-630 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 |
thomas-huber
changed the title
thomas-huber
marked this pull request as ready for review
August 27, 2026 06:41
Resolve conflicts in ce_coll.cc and enqueue.cc against develop's current shape, keeping develop's behavior for every path this work does not own. - Keep develop's inline CE plan construction in ncclLaunchPrepare rather than switching to scheduleCeCollTaskToPlan, which does not set useDda, ddaPeerBases, ddaUserRecvBuff, ddaCopyBackBytes, redOp or sizes and would regress CE AllReduce, the DDA scratch path and CE AlltoAllv. The hierarchical selection marker is added to develop's block instead. - Keep develop's node-count gate in ncclCeAvailable and ncclCeScratchAvailable. Hierarchical tasks already reach ceCollTaskAppend through hierCeAvailable, so relaxing the shared gate is unnecessary and would only affect scale-up domains that span a whole communicator. - Trigger CE init when hierarchical CE is selected. Develop moved the trigger into taskAppend behind nNodes == 1, so a multi-node communicator would otherwise never run ncclCeInit and would sync through a null pointer. - Size the CE sync window after ncclDevrInitOnce, which is what populates devrState.lsaSize. AICOMRCCL-630 Co-authored-by: Cursor <cursoragent@cursor.com>
…/AICOMRCCL-630-ce-scaleout Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/rccl/src/ce_coll.cc
Contributor
|
@thomas-huber - is the PR ready for review? This needs to go soon. |
Contributor
|
@AbandiGa - Can you review this PR? |
Contributor
|
@thomas-huber were you able to resolve the issue that was caused by the merge? Also, after resolving the issue can you test single-node CE as well. |
AbandiGa
reviewed
Sep 2, 2026
AbandiGa
approved these changes
Sep 2, 2026
nusislam
reviewed
Sep 2, 2026
Contributor
|
Unit tests add topology-gated MultiNodeHierarchical cases with three distinct epochs. They do not cover the 68 MiB / 64 MiB chunk boundary called out in the test plan. |
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Automated Review Guard for Upfront Scrutiny
Findings
1 finding(s) that fall outside the diff and could not be anchored to a line
projects/rccl/CHANGELOG.md:11(Low) No CHANGELOG entry for a user-visible, env-gated feature
Bottom line: Nothing new here blocks on top of the two gate bugs reviewers already filed at ce_coll.cc:567 and :2024. The two new hierarchical tests are the weak point: they gate on MPI topology only, so they false-red where RMA is absent.
Comment @argus review to re-run, or @argus re-review after pushing.
Correctness: - ncclMemOpSync: the graph-capture reset loop walked comm->nRanks while the ready/complete arrays and batchParams are both sized from the LSA team. With NCCL_CE_SYNC_OPS_PER_RANK_UC == 3 that overruns batchParams from four nodes up and indexes the flag arrays out of bounds from two nodes up. Walk lsaSize. - ncclLaunchCeColl: match the predicate ncclHierCeAvailable admits the task under (nNodes > 1 && !ncclDevrIsOneLsaTeam). Previously a single-node communicator with a reduced LSA team could be dispatched to the hierarchical path without its RMA prerequisites having been checked. - ncclCeInit: skip the CE AllReduce staging buffer unless the communicator is one LSA team. CE AllReduce is unreachable on the hierarchical path, so the allocation and its window registration were pure overhead there. Clarity: - Define RCCL_CE_HIER_SELECTED_TAG once and use it in both emitters and the test assertion, so the marker cannot drift from what the test matches. - Derive baseUCSymComplOffset from ceDevBaseSize. Tests: - Skip the scale-out cases when the communicator reports no symmetric memory support, so a stack without a GIN backend does not report a failure for behaving correctly. MPI topology alone does not imply the path is reachable. - Expect CE batch ops per LSA destination rather than per world rank. - Make the AlltoAll pattern depend on both endpoints so a wrong source slice changes the received bytes; the previous sender-only fill could not detect it. - Guard the LSA-scoped sync assertions with requireMinLsaRanks so two nodes at one rank each cannot pass them vacuously on an empty batch. - Pin the idempotence of ncclCeFreeBatchOpsParams, which hierarchical AllGather relies on when it releases per-chunk scratch and again at exit. - Cover the 64 MiB hierarchical chunk boundary with 68 MiB per rank and peer. AICOMRCCL-630 Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
There was a problem hiding this comment.
ARGUS
Findings
1 finding(s) that fall outside the diff and could not be anchored to a line
projects/rccl/src/ce_coll.cc:2029(Low) The PR description's Dispatch bullet no longer matches the shipped predicate
Bottom line: Worth settling before approval: scaleOutPrerequisitesMet() gates on one log string, so the documented NCCL_NUM_RMA_CTX=0 negative control reports a red on a correctly behaving stack.
Comment @argus review to re-run, or @argus re-review after pushing.
Make the hierarchical selection marker use the same predicate the launch dispatches on. ncclLaunchCeColl tests nNodes > 1 && !isOneLsaTeam while the enqueue.cc emitter tested only !isOneLsaTeam, so a single-node comm with a reduced LSA team printed the marker and then ran the LSA-local routine. The scale-out tests use that marker as their only oracle. Both sites now call a shared ncclCeHierDispatch. Reject CE for communicators whose LSA team does not cover them. The LSA-local routines address peers by LSA rank, so on a single node with NCCL_LSA_TEAM_SIZE below nRanks ncclCeAvailable admitted a task that would gather only lsaSize slices, at LSA rather than communicator indices. Such comms now fall back to the kernel path. Default behaviour is unchanged, since computeLsaSize yields nRanks on a single node when the variable is unset. Read the scale-out test prerequisites off the communicator instead of the log. The absence of "Symmetric memory is not supported" clears only the symmetricSupport clause of ncclHierCeAvailable, so a stack with GIN but no RMA context ran the case and failed on an absent marker rather than skipping. Predict intraBatchSync as false on the scale-out branch. Neither hierarchical collective assigns it, so the batch always logs the without-sync line; the 68 MiB chunk-boundary case already clears the byte threshold, so the prediction only held on nodes with at most kCeIntraBatchSyncFreq ranks. Cover the batch-params release fully: assert attrs, attrIdxs, numAttrs and intraBatchSync, and assert the pointers are non-null after init so the nulling checks cannot pass vacuously. Add host coverage for ncclDevrGetWinOffset in DevRuntimeTests, which already includes the hipified dev_runtime.cc, so no new source seam is needed. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
…/AICOMRCCL-630-ce-scaleout Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/rccl/CHANGELOG.md
Contributor
There was a problem hiding this comment.
ARGUS
Findings
1 finding(s) that fall outside the diff and could not be anchored to a line
projects/rccl/src/rccl_wrap.cc:1173(Low) Enabling multi-node CE makes rcclSelectAllGather report the wrong algorithm, since its CE branch tests only ncclCeAvailable
Bottom line: Worth settling before approval: the CE-init trigger at enqueue.cc:3882 now keys on a predicate that varies by node, so a job with uneven ranks per node joins the symmetric-register group on only some ranks.
Comment @argus review to re-run, or @argus re-review after pushing.
rcclSelectAllGather only tested ncclCeAvailable, so multi-node CE AllGather was labeled Hier/Direct/RING even when taskAppend took hierarchical CE.
Mocks now skip computeLsaSize's null rankToNode walk, maxLocalRanks is set on the hier helper, and FreeBatchOpsParamsIsIdempotent is in CE_Internal_Init.
…/AICOMRCCL-630-ce-scaleout Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/rccl/CHANGELOG.md
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: the new !ncclDevrIsOneLsaTeam reject in ncclCeScratchAvailable has no negative case, so deleting it leaves the eligibility suite green.
Comment @argus review to re-run, or @argus re-review after pushing.
mustafabar
self-requested a review
September 16, 2026 15:06
mustafabar
reviewed
Sep 16, 2026
Contributor
There was a problem hiding this comment.
How does the heatmap performance scale on MI300 and MI350 systems at 4-node and 8-node configurations? Given that these changes affect a large number of source files on the critical path of execution, and although the feature is disabled by default, it is difficult to classify the change set as low regression risk since they are compiled in.
Additionally, what is the minimum set of environment variables required to enable H CE-scaleout? The description lists several settings:
NCCL_CTA_POLICY=2
NCCL_LOCAL_REGISTER=0
NCCL_CUMEM_ENABLE=1
NCCL_GIN_ENABLE=1
NCCL_NUM_RMA_CTX>0
Are all of these required, or are some already satisfied by default?
For the minimum required set, what performance improvement is observed compared to the baseline?
AbandiGa
approved these changes
Sep 16, 2026
Add negative cases for the !ncclDevrIsOneLsaTeam reject in ncclCeScratchAvailable, which had no coverage: every mock comm in the CE eligibility suites has lsaSize == nRanks, so deleting the clause left both suites green. The new cases use lsaSize = nRanks/2 with nNodes = 1 and were confirmed to fail with the clause removed. Rewrite the LSA-coverage CHANGELOG entry to lead with the symptom, silently incorrect data under NCCL_LSA_TEAM_SIZE, rather than the internal condition. Correct the enablement entry as well: NCCL_GIN_ENABLE, NCCL_NUM_RMA_CTX, NCCL_WIN_ENABLE and NCCL_DMABUF_ENABLE already default to qualifying values, NCCL_LOCAL_REGISTER is not read by any selection gate, and NCCL_CUMEM_ENABLE=1 is required on gfx942/gfx950 because cuMem auto-detect is limited to gfx1250. Co-authored-by: Cursor <cursoragent@cursor.com>
garzaran
approved these changes
Sep 16, 2026
garzaran
left a comment
garzaran
left a comment
There was a problem hiding this comment.
This looks good. I know we are still doing performance validation, but the changes seem correct.
Contributor
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: with NCCL_RMA_DISABLE=1 a hierarchical CE task is still admitted, and ncclLaunchCeColl dispatches it into a rmaProxyCtxs array that was never allocated.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
|
@thomas-huber - there are several conflicts. You should fix these and merge the PR. |
…/AICOMRCCL-630-ce-scaleout Absorbs the NCCL v2.31.2-1 sync (#10785). Develop independently landed most of this branch's CE-side fixes during sync stabilization, so the resolution keeps develop's version of ce_coll.cc, dev_runtime.cc and rma_proxy_launch.cc and re-applies only what is still missing. Carries one behaviour fix of its own, in rma_proxy.cc: ncclRmaWantInternalCtx now defers to ncclRmaProxyEnabled instead of re-deriving its terms. This is the amd-argus finding on comment 4029442606, which survived the sync. ncclHierCeAvailable gates on ncclRmaWantInternalCtx, but the contexts it promises are allocated by ncclRmaProxyConnectOnce, which only runs when ncclRmaProxyEnabled holds. NCCL_RMA_DISABLE=1 cleared one side and not the other, so a multi-node comm with NCCL_CTA_POLICY=2 and symmetric send and recv windows admitted the hierarchical path and then dereferenced a NULL rmaProxyCtxs at launch. Composing the predicates also picks up numRmaCtx > 0 and globalRmaProxySupport, so ncclHierCeAvailable now provably implies ncclRmaProxyEnabled, and it fixes the numIntCtx line that shares the predicate. Superseded by develop, dropped here: - LSA-indexed ready/complete slots in ncclPrepUCSync* and ncclMemOpSync. Develop iterates the LSA team but keeps the slots indexed by global rank via devrState.lsaRankList, which is self-consistent; keeping our LSA-indexed writer against develop's global-rank waiter would have deadlocked. - The isOneLsaTeam guard on the CE AllReduce staging buffer. Develop allocates it lazily in ncclCeEnsureAllReduceStaging, which covers more cases. - ncclDevrGetRmaWin returning rmaHostWins, ncclDevrGetWinOffset and the ncclRmaProxyPutBuildOp offsets. Develop carries all three and additionally indexes data-window handles by rmaProxyCtx->collCommIdx. - The idempotent ncclCeFreeBatchOpsParams and the ncclCeScratchAvailable split. - useMCSync's !p2pCrossClique term, now expressed as symkState.hasLsaMultimem. Adapted to develop's new semantics: - ncclCeAvailable and ncclHierCeAvailable take sendWin/recvWin and reject host-backed cuMem segments. The ncclHierCeAvailable host fake had to follow: nothing asserts drift on that hook, so rccl-UnitTestsMicroEnqueue would have failed to link against the real 7-argument declaration. - The hierarchical rail draws on the internal RMA context range [numRmaCtx, numRmaCtx + numIntCtx) rather than the user-addressable contexts, so NCCL_NUM_RMA_CTX no longer sizes it. - ncclMemOpSync selects MC vs UC sync off a single useMCSync local so the op-count sizing cannot under-allocate the path it then takes. - CeAlltoAllvTestHelpers sets globalRmaProxySupport. init.cc derives hostRmaSupport from it on a multi-clique comm and ncclRmaProxyEnabled reads it directly, so a mock that set only hostRmaSupport described a comm that cannot exist. Still ours, re-applied: - ncclHierCeDispatch, shared by the launch site and the selection marker. - maxLocalRanks rather than localRanks in the LSA-coverage clause, so the predicate is comm-uniform and cannot split a group join. - The CE-init trigger widened to reach multi-node hierarchical communicators. - LSA coverage on ncclCeScratchAvailable, which develop left only on ncclCeAvailable, so the RCCL_FORCE_CE scratch path was still exposed. - Hierarchical CE AllGather reported as CE for -A 1. Co-authored-by: Cursor <cursoragent@cursor.com>
…t gate ncclCeFreeBatchOpsParams came out of the NCCL v2.31.2-1 sync keeping the pointer nulling but not the counter resets, so numOps, intraBatchSync and numAttrs survived a release. The hierarchical collectives free the same params twice, once per chunk batch and once at exit, and a stale numOps makes the second release look like it still describes a populated batch. CeInternalMPITest.FreeBatchOpsParamsIsIdempotent covers exactly this and was failing; it passes again. Adds CeAlltoAllEligibilityTest.WantInternalCtxRequiresProxyEnabled for the ncclRmaWantInternalCtx change in the merge commit. It pins the composition through numRmaCtx, the one term that appears in ncclRmaProxyEnabled and in no clause of ncclRmaWantInternalCtx's own: set it to 0 on an otherwise hier-eligible comm and both ncclRmaWantInternalCtx and ncclHierCeAvailable must decline. Re-deriving the terms locally, in any form that omits the deferral, makes this case fail. The NCCL_RMA_DISABLE path is deliberately not driven through the environment. NCCL_PARAM caches on first read, so such a case would pass or fail depending on whether anything in the binary read the variable first, and an order-dependent test is worse than none. The composition test kills the same mutant without that fragility. The remaining edits in this file are the sendWin/recvWin arguments the sync added to ncclCeAvailable and ncclHierCeAvailable. Co-authored-by: Cursor <cursoragent@cursor.com>
…e-out Neither of these is CE scale-out's, and both can be dropped from this PR if you would rather see them separately. They are here because this branch cannot be verified without them: rccl-UnitTestsDevRuntime segfaults on the first, and the --debug --tests_build --enable-mpi-tests build does not compile at all with the second, and that is the configuration this branch's CE suites run in. 1. symMemoryDestroy lost its membership check (dev_runtime.cc). Before the sync it located mem in devr->memHead and returned early when it was absent, which is what made the function safe to call twice. The guard and the test that covers it both came from a7ad311, "harden symmetric GIN teardown on gfx1250 VMM paths". The post-sync version does the whole teardown first and only then unlinks, with an unguarded walk: struct ncclDevrMemory** ptr = &devr->memHead; while (*ptr != mem) ptr = &(*ptr)->next; This is reachable in production, not just from the test. The error ladders in ncclDevrWindowRegisterInGroup and ncclDevrCommCreateInternal fall through from symWindowDestroy, which destroys winHost->memory, into symMemoryDestroy(comm, mem) on the same memory. A failed symmetric window registration therefore repeats every ncclGinDeregister, ncclRmaProxyDeregister, cuMemUnmap, cuMemRelease and free, then runs off the end of the list. It affects every symmetric-window user, GIN and the Device API included. CI does not see it because CI builds release without --tests_build and nothing runs ctest. Develop's own SymMemoryObtainTest.DestroyIsIdempotent segfaults on it. The check is back at the top, before anything is touched. Note the drain loop at dev_runtime.cc:330 passes devr->memHead, which is always a member, so the early return cannot turn that loop into a spin. 2. GdrFlushTest.ForcedScratchpadWrite_ReproducesFault does not compile. The sync added the case but not the two things it calls: RunRecvFlushBurst has no forceWrite parameter and there is no gdrSupported(). Guarded off rather than half-fixed, because restoring the forceWrite hook is a decision for whoever owns the GDR flush fence. It only breaks builds with ENABLE_FAULT_INJECTION on, which is the default, combined with test targets, which CI omits. Found but not fixed here: SkipCuMemFreePolicy.IsolatedArchAndEnvBranches fails on pristine develop's dev_runtime.cc as well, verified by reverting this file and rebuilding. It concerns cuMem free policy and is nothing to do with CE, so rccl-UnitTestsDevRuntime is expected to read 5/6 rather than 6/6. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
@argus re-review Head is now Three things worth pointing you at:
|
…equency
CollImplInfo.AllReduceMatchesDispatchLog dies with SIGFPE on gfx94X. The graph-
captured unicast branch of the CE cost model divides by
comm->ceColl.intraBatchSyncFreq:
int freq = (int)comm->ceColl.intraBatchSyncFreq;
if (numOps > freq && (uint64_t)perRankBytes * (uint64_t)numOps >= comm->ceColl.intraBatchSyncMsgThreshold) {
int blocks = (numOps + freq - 1) / freq;
That field and intraBatchSyncMsgThreshold are written in exactly one place,
ncclCeInit (ce_coll.cc:197-198), so both are still 0 on a communicator that
never initialized CE. gfx942 is such a communicator: ncclCeImplemented declines
on the driver version, so ncclCeInit never runs.
With freq == 0 the existing condition does not protect the divide, it
guarantees it. numOps > 0 holds for any rank count, and the threshold term
compares against 0, so both halves pass and the divide is reached every time.
gfx950 hides this because CE does initialize there and freq is non-zero.
Adding freq > 0 is the whole fix: with no CE init there is no intra-batch sync,
so there is no latency for the model to add and skipping the term is correct.
src/tuning/ce_model.cc arrived with the NCCL v2.31.2-1 sync (#10785) and does
not exist at this branch's pre-sync base, and the file is otherwise identical to
develop, so this is a sync regression rather than CE scale-out's. It is fixed
here because it fails a required check on this PR.
Verified: CollImplInfo.* is 9/9 on gfx950 before and after, so the guard changes
nothing where freq was already valid.
Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
CI still pending checks at 10PM my time, hesitant to merge without full green checks. cc @nusislam |
Contributor
Author
Hey @mustafabar the heatmap performance scaling on 4 and 8 nodes came back totally clean for MI300 and MI350 (8 GPUS 8NICS per node). Per the env vars, only two of the five are explicitly required on Ruby for example. NCCL_CTA_POLICY=2 is required. NCCL_CUMEM_ENABLE=1 is required. The other three are already satisfied: NCCL_GIN_ENABLE defaults to 1, NCCL_NUM_RMA_CTX defaults to 1 so >0 holds out of the box, and NCCL_DMABUF_ENABLE gets force-enabled automatically once cuMem is on. NCCL_LOCAL_REGISTER=0 isn't in the gate chain at all. I will update the description to reflect this. |
thomas-huber
deleted the
users/thomas-huber/AICOMRCCL-630-ce-scaleout
branch
September 18, 2026 15:09
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task
amd-jiali
added a commit
that referenced
this pull request
Sep 29, 2026
## Motivation - The NCCL v2.31.2-1 sync (`b71e057a71c`) added `GdrFlushTest.ForcedScratchpadWrite_ReproducesFault` under `ENABLE_FAULT_INJECTION`. That case calls `gdrSupported()` and `RunRecvFlushBurst(..., /*forceWrite=*/true, &forced)`. There is no `gdrSupported` symbol, and `RunRecvFlushBurst` takes three arguments. `FAULT_INJECTION` defaults to `ON`, so `rccl-UnitTestsMPI` fails to compile. The nightly then reports return code 183 because OpenMPI cannot start the missing binary. The coco PR arm does not build this target, so the break only shows up in `oci-rccl-unit-nightly`. - `#10450` wrapped the case in `#if 0`. The calls were still the merge-introduced ones, and the case cannot be restored without a `forceWrite` hook that was never added. ## Technical Details Remove `ForcedScratchpadWrite_ReproducesFault` from `projects/rccl/test/transport/NetIbMPI/GdrFlushTests.cpp`. It was not in the tree before the sync. The other GDR flush cases keep the three-argument `RunRecvFlushBurst` and `gdrPtrSupport()`. ## Issue Tracking JIRA ID: AICOMRCCL-2403 ## Test Plan - On a compute node with ROCm 7.14.0, from a tree that contains `b71e057a71c`, rebuild `rccl-UnitTestsMPI` with `ENABLE_MPI_TESTS=ON` and `FAULT_INJECTION=ON`. - Confirm `GdrFlushTests.cpp` compiles, the binary is produced, and `ForcedScratchpadWrite` is not linked. ## Test Result The debug cache had `FAULT_INJECTION=ON` and `ENABLE_MPI_TESTS=ON`. `GdrFlushTests.cpp` compiled and `rccl-UnitTestsMPI` linked. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
amd-jiali
added a commit
that referenced
this pull request
Sep 30, 2026
…e drift (#12054) ## Motivation The NCCL v2.31.2-1 sync added `sendWin` and `recvWin` to `ncclHierCeAvailable` and `ncclCeAvailable`. The debug gtest stack failed to compile because `CeAlltoAllEligibilityTest.MultiNodeHierAvailable_DoesNotYieldDda` was still the 5-argument form. #10450 updated those calls, and the host fake is already the 7-parameter form on `develop`. Two gaps from that same change were still open: - `ncclHierCeAvailable` in `test/host/fakes/ce_fakes.cc` had no `ASSERT_HOOK_MATCHES_PROD` seam, so a later signature change can compile as a silent extra overload. - The debug eligibility tests compile and then fail. `configureHierEligible()` never sets `hierarchicalCommsInitialized`, which `ncclHierCeAvailable` now requires. The LSA-team scratch test still passes `ncclSymSendRegRecvReg`, and `ncclCeScratchAvailable` now rejects a registered receive window. ## Technical Details - `ncclHierCeAvailable` now goes through `g_hierCeAvailableFn`. `ASSERT_HOOK_MATCHES_PROD(g_hierCeAvailableFn, ncclHierCeAvailable)` matches the existing checks for `ncclCeAvailable`. Taking `&ncclHierCeAvailable` makes a leftover extra overload ill-formed. `g_hierCeAvailable` stays the default return value. - `CeAlltoAllvMockComm::configureHierEligible()` sets `hierarchicalCommsInitialized`. Init only sets that flag when `nNodes >= 8`. `UnequalRanksPerNode_HierUnavailable` and `WantInternalCtxRequiresProxyEnabled` still reject on their own gates. - `CeAvailable_LsaTeamSmallerThanCommRejected` calls `ncclCeScratchAvailable` with `ncclSymSendRegRecvNonreg`. Restoring `lsaSize` flips both `ncclCeScratchAvailable` and `ncclCeAvailable`. The registered-window case stays on `ncclCeAvailable`. ## Issue Tracking JIRA ID: AICOMRCCL-2402 ## Test Plan On a compute node with ROCm 7.14.0 (`gfx950`): - `./install.sh --debug -l -t` builds `rccl-UnitTestsFixturesDebug` and compiles `test/host/fakes/ce_fakes.cc` into `rccl-UnitTestsMicro`. - Run: `rccl-UnitTestsFixturesDebug --gtest_filter='CeAlltoAllEligibilityTest.MultiNodeHierAvailable*:CeAlltoAllEligibilityTest.UnequalRanksPerNode*:CeAlltoAllEligibilityTest.WantInternalCtx*:CeAlltoAllvEligibilityTest.CeAvailable*'` ## Test Result - The debug stack built. No `no matching function for call to 'ncclHierCeAvailable'` or `'ncclCeAvailable`'. ce_fakes.cc compiled into rccl-UnitTestsMicro. - After the fixture update: the same filter, 8 tests, all passed. A one-parameter change to the hier declaration failed `ce_fakes.cc` at compile time with signature drift: `g_hierCeAvailableFn` no longer matches `ncclHierCeAvailable`. The header was restored. The matching 7-parameter declaration compiles. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
maple52046
pushed a commit
to maple52046/rocm-systems
that referenced
this pull request
Oct 1, 2026
…ifts The debug gtest AlltoAll eligibility calls were updated with sendWin/recvWin in ROCm#10450, but the hier fake had no ASSERT_HOOK_MATCHES_PROD seam, so a later signature change could again compile as a silent extra overload. Co-authored-by: Cursor <cursoragent@cursor.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
## Motivation CE collectives in RCCL operate within a single LSA team, i.e. one node. The NCCL 2.30.7 sync brought the hierarchical scale-out CE implementation into develop; and this change completes its integration into the launch, scheduling and synchronization paths. With this PR, `ncclAllGather` and `ncclAlltoAll` run hierarchically on multi-node communicators: intra-node over SDMA / Copy Engine across the LSA team, inter-node over one-sided RMA through the CPU (GIN) proxy. ## Technical Details - **Dispatch** (`src/ce_coll.cc`): `ncclLaunchCeColl` now routes AllGather and AllToAll to `ncclHierCeAllGather` / `ncclHierCeAlltoAll` when the communicator spans more than one LSA team, and to the existing LSA-local routines otherwise. The predicate is `!ncclDevrIsOneLsaTeam(comm)` rather than `comm->nNodes`, so a multi-clique single-NVLink-domain communicator still takes the LSA-local path. Functions without a hierarchical implementation return `ncclInvalidUsage`. - **Scheduling** (`src/enqueue.cc`): `ncclLaunchPrepare` now calls `scheduleCeCollTaskToPlan`, replacing the inline CE plan construction. - **LSA-scoped synchronization** (`src/ce_coll.cc`): CE sync state and barriers are sized and indexed by the LSA team rather than by the world communicator, since the LSA pointer helpers address LSA-local ranks. `ncclCeInit`, `ncclPrepMCSync`, `ncclPrepUCSync` and `ncclMemOpSync` use `devrState.lsaSize` and `devrState.lsaSelf`. - **Availability gate**: `ncclCeAvailable` tests LSA coverage of the communicator via `ncclTeamLsa(comm).nRanks < comm->nRanks`, matching upstream, in place of a node-count test. - **Multicast synchronization** is not available across scale-up cliques, so `ncclMemOpSync` selects it with `nvlsSupport && !p2pCrossClique` and falls back to unicast otherwise. - **RMA proxy handles** (`src/dev_runtime.cc`): `ncclDevrGetRmaWin` returns the host RMA handles produced by `ncclRmaProxyRegister`, which is what the host proxy path consumes for data-bearing operations. Adds `ncclDevrGetWinOffset` and applies it to `srcOff` / `dstOff` in `ncclRmaProxyPutBuildOp`, porting NCCL `1e47331c24ecc1a2e42ace9a0932f100919ba1ad` so that a window registered within a larger backing allocation resolves to the correct address. - **Batch cleanup**: `ncclCeFreeBatchOpsParams` clears the pointers and counters it releases, so the per-chunk and end-of-function cleanup paths in hierarchical AllGather compose safely. Tests add `CeMPI_AllGather.MultiNodeHierarchical` and `CeMPI_AlltoAll.MultiNodeHierarchical`, topology-gated via `MPI_Comm_split_type` so they skip rather than pass vacuously on a single node. Each runs three data-distinct iterations on one communicator, so a stale buffer cannot satisfy the check, and asserts the `[Hierarchical CE]` selection marker alongside payload correctness. The white-box `ncclPrepUCSync` expectations in `CeInternalMPITests.cpp` move to the LSA-local rank domain to match. Selecting this path requires `NCCL_CTA_POLICY=2`, `NCCL_CUMEM_ENABLE=1`, `NCCL_GIN_ENABLE=1`, `NCCL_NUM_RMA_CTX>0` and symmetric windows. The GIN requirement follows upstream 2.30.7: `ncclHierCeAvailable` is unchanged from `v2.30.7-1` and `init.cc` is untouched by this PR. There is no message-size term in the availability check; payloads above the 64 MiB per rank/peer chunk are chunked rather than falling back. ## Issue Tracking JIRA ID: AICOMRCCL-630 Graph-capture CE synchronization is tracked separately under AICOMRCCL-1548 and is not included here. ## Test Plan MI350X, 8 GPUs and 8 ranks per node, RoCE fabric, ROCm 7.14. Scales of 1, 2, 4, 8, 16 and 32 nodes. - All existing single-node CE cases, plus the two new multi-node cases run as two fresh-communicator cycles per scale. - rccl-tests AllGather and AllToAll sweeps, 4 KiB through 1 GiB. - Negative controls: `NCCL_CTA_POLICY=0`, and `NCCL_NUM_RMA_CTX=0` with symmetric windows — correct results with the hierarchical path not selected. - 68 MiB per rank/peer, crossing the 64 MiB chunk boundary. - A/B control toggling `NCCL_GIN_ENABLE`, confirming the symmetric-memory requirement empirically. - CE vs default vs `RCCL_HIERARCHICAL_ALLGATHER` for AllGather, and CE vs default for AllToAll, at 4/8/16/32 nodes on a release build. Also tested MI350 and MI300 4 and 8 nodes to compare develop vs develop + my branch to ensure there's no perf regression, sweeps came back clean. Each arm is gated on the named gtest reporting `[ OK ]` and not SKIPPED, data correctness, the `[Hierarchical CE]` marker, RMA proxy assignment on every rank, and absence of timeouts or HIP errors. ## Test Result **Correctness passes at every scale from 1 to 32 nodes (256 GPUs), with no skipped tests.** Representative 32-node CE AllGather: `Out of bounds values : 0 OK`, hierarchical selection markers present, RMA proxy assigned on all 256 ranks, 2560 RMA completions, no warnings. See ticket for performance impacts. ## 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>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
## Motivation CE collectives in RCCL operate within a single LSA team, i.e. one node. The NCCL 2.30.7 sync brought the hierarchical scale-out CE implementation into develop; and this change completes its integration into the launch, scheduling and synchronization paths. With this PR, `ncclAllGather` and `ncclAlltoAll` run hierarchically on multi-node communicators: intra-node over SDMA / Copy Engine across the LSA team, inter-node over one-sided RMA through the CPU (GIN) proxy. ## Technical Details - **Dispatch** (`src/ce_coll.cc`): `ncclLaunchCeColl` now routes AllGather and AllToAll to `ncclHierCeAllGather` / `ncclHierCeAlltoAll` when the communicator spans more than one LSA team, and to the existing LSA-local routines otherwise. The predicate is `!ncclDevrIsOneLsaTeam(comm)` rather than `comm->nNodes`, so a multi-clique single-NVLink-domain communicator still takes the LSA-local path. Functions without a hierarchical implementation return `ncclInvalidUsage`. - **Scheduling** (`src/enqueue.cc`): `ncclLaunchPrepare` now calls `scheduleCeCollTaskToPlan`, replacing the inline CE plan construction. - **LSA-scoped synchronization** (`src/ce_coll.cc`): CE sync state and barriers are sized and indexed by the LSA team rather than by the world communicator, since the LSA pointer helpers address LSA-local ranks. `ncclCeInit`, `ncclPrepMCSync`, `ncclPrepUCSync` and `ncclMemOpSync` use `devrState.lsaSize` and `devrState.lsaSelf`. - **Availability gate**: `ncclCeAvailable` tests LSA coverage of the communicator via `ncclTeamLsa(comm).nRanks < comm->nRanks`, matching upstream, in place of a node-count test. - **Multicast synchronization** is not available across scale-up cliques, so `ncclMemOpSync` selects it with `nvlsSupport && !p2pCrossClique` and falls back to unicast otherwise. - **RMA proxy handles** (`src/dev_runtime.cc`): `ncclDevrGetRmaWin` returns the host RMA handles produced by `ncclRmaProxyRegister`, which is what the host proxy path consumes for data-bearing operations. Adds `ncclDevrGetWinOffset` and applies it to `srcOff` / `dstOff` in `ncclRmaProxyPutBuildOp`, porting NCCL `1e47331c24ecc1a2e42ace9a0932f100919ba1ad` so that a window registered within a larger backing allocation resolves to the correct address. - **Batch cleanup**: `ncclCeFreeBatchOpsParams` clears the pointers and counters it releases, so the per-chunk and end-of-function cleanup paths in hierarchical AllGather compose safely. Tests add `CeMPI_AllGather.MultiNodeHierarchical` and `CeMPI_AlltoAll.MultiNodeHierarchical`, topology-gated via `MPI_Comm_split_type` so they skip rather than pass vacuously on a single node. Each runs three data-distinct iterations on one communicator, so a stale buffer cannot satisfy the check, and asserts the `[Hierarchical CE]` selection marker alongside payload correctness. The white-box `ncclPrepUCSync` expectations in `CeInternalMPITests.cpp` move to the LSA-local rank domain to match. Selecting this path requires `NCCL_CTA_POLICY=2`, `NCCL_CUMEM_ENABLE=1`, `NCCL_GIN_ENABLE=1`, `NCCL_NUM_RMA_CTX>0` and symmetric windows. The GIN requirement follows upstream 2.30.7: `ncclHierCeAvailable` is unchanged from `v2.30.7-1` and `init.cc` is untouched by this PR. There is no message-size term in the availability check; payloads above the 64 MiB per rank/peer chunk are chunked rather than falling back. ## Issue Tracking JIRA ID: AICOMRCCL-630 Graph-capture CE synchronization is tracked separately under AICOMRCCL-1548 and is not included here. ## Test Plan MI350X, 8 GPUs and 8 ranks per node, RoCE fabric, ROCm 7.14. Scales of 1, 2, 4, 8, 16 and 32 nodes. - All existing single-node CE cases, plus the two new multi-node cases run as two fresh-communicator cycles per scale. - rccl-tests AllGather and AllToAll sweeps, 4 KiB through 1 GiB. - Negative controls: `NCCL_CTA_POLICY=0`, and `NCCL_NUM_RMA_CTX=0` with symmetric windows — correct results with the hierarchical path not selected. - 68 MiB per rank/peer, crossing the 64 MiB chunk boundary. - A/B control toggling `NCCL_GIN_ENABLE`, confirming the symmetric-memory requirement empirically. - CE vs default vs `RCCL_HIERARCHICAL_ALLGATHER` for AllGather, and CE vs default for AllToAll, at 4/8/16/32 nodes on a release build. Also tested MI350 and MI300 4 and 8 nodes to compare develop vs develop + my branch to ensure there's no perf regression, sweeps came back clean. Each arm is gated on the named gtest reporting `[ OK ]` and not SKIPPED, data correctness, the `[Hierarchical CE]` marker, RMA proxy assignment on every rank, and absence of timeouts or HIP errors. ## Test Result **Correctness passes at every scale from 1 to 32 nodes (256 GPUs), with no skipped tests.** Representative 32-node CE AllGather: `Out of bounds values : 0 OK`, hierarchical selection markers present, RMA proxy assigned on all 256 ranks, 2560 RMA completions, no warnings. See ticket for performance impacts. ## 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>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
## Motivation CE collectives in RCCL operate within a single LSA team, i.e. one node. The NCCL 2.30.7 sync brought the hierarchical scale-out CE implementation into develop; and this change completes its integration into the launch, scheduling and synchronization paths. With this PR, `ncclAllGather` and `ncclAlltoAll` run hierarchically on multi-node communicators: intra-node over SDMA / Copy Engine across the LSA team, inter-node over one-sided RMA through the CPU (GIN) proxy. ## Technical Details - **Dispatch** (`src/ce_coll.cc`): `ncclLaunchCeColl` now routes AllGather and AllToAll to `ncclHierCeAllGather` / `ncclHierCeAlltoAll` when the communicator spans more than one LSA team, and to the existing LSA-local routines otherwise. The predicate is `!ncclDevrIsOneLsaTeam(comm)` rather than `comm->nNodes`, so a multi-clique single-NVLink-domain communicator still takes the LSA-local path. Functions without a hierarchical implementation return `ncclInvalidUsage`. - **Scheduling** (`src/enqueue.cc`): `ncclLaunchPrepare` now calls `scheduleCeCollTaskToPlan`, replacing the inline CE plan construction. - **LSA-scoped synchronization** (`src/ce_coll.cc`): CE sync state and barriers are sized and indexed by the LSA team rather than by the world communicator, since the LSA pointer helpers address LSA-local ranks. `ncclCeInit`, `ncclPrepMCSync`, `ncclPrepUCSync` and `ncclMemOpSync` use `devrState.lsaSize` and `devrState.lsaSelf`. - **Availability gate**: `ncclCeAvailable` tests LSA coverage of the communicator via `ncclTeamLsa(comm).nRanks < comm->nRanks`, matching upstream, in place of a node-count test. - **Multicast synchronization** is not available across scale-up cliques, so `ncclMemOpSync` selects it with `nvlsSupport && !p2pCrossClique` and falls back to unicast otherwise. - **RMA proxy handles** (`src/dev_runtime.cc`): `ncclDevrGetRmaWin` returns the host RMA handles produced by `ncclRmaProxyRegister`, which is what the host proxy path consumes for data-bearing operations. Adds `ncclDevrGetWinOffset` and applies it to `srcOff` / `dstOff` in `ncclRmaProxyPutBuildOp`, porting NCCL `1e47331c24ecc1a2e42ace9a0932f100919ba1ad` so that a window registered within a larger backing allocation resolves to the correct address. - **Batch cleanup**: `ncclCeFreeBatchOpsParams` clears the pointers and counters it releases, so the per-chunk and end-of-function cleanup paths in hierarchical AllGather compose safely. Tests add `CeMPI_AllGather.MultiNodeHierarchical` and `CeMPI_AlltoAll.MultiNodeHierarchical`, topology-gated via `MPI_Comm_split_type` so they skip rather than pass vacuously on a single node. Each runs three data-distinct iterations on one communicator, so a stale buffer cannot satisfy the check, and asserts the `[Hierarchical CE]` selection marker alongside payload correctness. The white-box `ncclPrepUCSync` expectations in `CeInternalMPITests.cpp` move to the LSA-local rank domain to match. Selecting this path requires `NCCL_CTA_POLICY=2`, `NCCL_CUMEM_ENABLE=1`, `NCCL_GIN_ENABLE=1`, `NCCL_NUM_RMA_CTX>0` and symmetric windows. The GIN requirement follows upstream 2.30.7: `ncclHierCeAvailable` is unchanged from `v2.30.7-1` and `init.cc` is untouched by this PR. There is no message-size term in the availability check; payloads above the 64 MiB per rank/peer chunk are chunked rather than falling back. ## Issue Tracking JIRA ID: AICOMRCCL-630 Graph-capture CE synchronization is tracked separately under AICOMRCCL-1548 and is not included here. ## Test Plan MI350X, 8 GPUs and 8 ranks per node, RoCE fabric, ROCm 7.14. Scales of 1, 2, 4, 8, 16 and 32 nodes. - All existing single-node CE cases, plus the two new multi-node cases run as two fresh-communicator cycles per scale. - rccl-tests AllGather and AllToAll sweeps, 4 KiB through 1 GiB. - Negative controls: `NCCL_CTA_POLICY=0`, and `NCCL_NUM_RMA_CTX=0` with symmetric windows — correct results with the hierarchical path not selected. - 68 MiB per rank/peer, crossing the 64 MiB chunk boundary. - A/B control toggling `NCCL_GIN_ENABLE`, confirming the symmetric-memory requirement empirically. - CE vs default vs `RCCL_HIERARCHICAL_ALLGATHER` for AllGather, and CE vs default for AllToAll, at 4/8/16/32 nodes on a release build. Also tested MI350 and MI300 4 and 8 nodes to compare develop vs develop + my branch to ensure there's no perf regression, sweeps came back clean. Each arm is gated on the named gtest reporting `[ OK ]` and not SKIPPED, data correctness, the `[Hierarchical CE]` marker, RMA proxy assignment on every rank, and absence of timeouts or HIP errors. ## Test Result **Correctness passes at every scale from 1 to 32 nodes (256 GPUs), with no skipped tests.** Representative 32-node CE AllGather: `Out of bounds values : 0 OK`, hierarchical selection markers present, RMA proxy assigned on all 256 ranks, 2560 RMA completions, no warnings. See ticket for performance impacts. ## 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>
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
CE collectives in RCCL operate within a single LSA team, i.e. one node. The
NCCL 2.30.7 sync brought the hierarchical scale-out CE implementation into
develop; and this change completes its integration into the launch,
scheduling and synchronization paths.
With this PR,
ncclAllGatherandncclAlltoAllrun hierarchically onmulti-node communicators: intra-node over SDMA / Copy Engine across the LSA
team, inter-node over one-sided RMA through the CPU (GIN) proxy.
Technical Details
src/ce_coll.cc):ncclLaunchCeCollnow routes AllGather andAllToAll to
ncclHierCeAllGather/ncclHierCeAlltoAllwhen thecommunicator spans more than one LSA team, and to the existing LSA-local
routines otherwise. The predicate is
!ncclDevrIsOneLsaTeam(comm)ratherthan
comm->nNodes, so a multi-clique single-NVLink-domain communicatorstill takes the LSA-local path. Functions without a hierarchical
implementation return
ncclInvalidUsage.src/enqueue.cc):ncclLaunchPreparenow callsscheduleCeCollTaskToPlan, replacing the inline CE plan construction.src/ce_coll.cc): CE sync state and barriersare sized and indexed by the LSA team rather than by the world communicator,
since the LSA pointer helpers address LSA-local ranks.
ncclCeInit,ncclPrepMCSync,ncclPrepUCSyncandncclMemOpSyncusedevrState.lsaSizeanddevrState.lsaSelf.ncclCeAvailabletests LSA coverage of thecommunicator via
ncclTeamLsa(comm).nRanks < comm->nRanks, matchingupstream, in place of a node-count test.
ncclMemOpSyncselects it withnvlsSupport && !p2pCrossCliqueand fallsback to unicast otherwise.
src/dev_runtime.cc):ncclDevrGetRmaWinreturns thehost RMA handles produced by
ncclRmaProxyRegister, which is what the hostproxy path consumes for data-bearing operations. Adds
ncclDevrGetWinOffsetand applies it to
srcOff/dstOffinncclRmaProxyPutBuildOp, portingNCCL
1e47331c24ecc1a2e42ace9a0932f100919ba1adso that a window registeredwithin a larger backing allocation resolves to the correct address.
ncclCeFreeBatchOpsParamsclears the pointers and countersit releases, so the per-chunk and end-of-function cleanup paths in
hierarchical AllGather compose safely.
Tests add
CeMPI_AllGather.MultiNodeHierarchicalandCeMPI_AlltoAll.MultiNodeHierarchical, topology-gated viaMPI_Comm_split_typeso they skip rather than pass vacuously on a single node.Each runs three data-distinct iterations on one communicator, so a stale buffer
cannot satisfy the check, and asserts the
[Hierarchical CE]selection markeralongside payload correctness. The white-box
ncclPrepUCSyncexpectations inCeInternalMPITests.cppmove to the LSA-local rank domain to match.Selecting this path requires
NCCL_CTA_POLICY=2,NCCL_CUMEM_ENABLE=1,NCCL_GIN_ENABLE=1,NCCL_NUM_RMA_CTX>0and symmetricwindows. The GIN requirement follows upstream 2.30.7:
ncclHierCeAvailableis unchanged fromv2.30.7-1andinit.ccis untouchedby this PR. There is no message-size term in the availability check; payloads
above the 64 MiB per rank/peer chunk are chunked rather than falling back.
Issue Tracking
JIRA ID: AICOMRCCL-630
Graph-capture CE synchronization is tracked separately under AICOMRCCL-1548 and
is not included here.
Test Plan
MI350X, 8 GPUs and 8 ranks per node, RoCE fabric, ROCm 7.14. Scales of 1, 2, 4,
8, 16 and 32 nodes.
two fresh-communicator cycles per scale.
NCCL_CTA_POLICY=0, andNCCL_NUM_RMA_CTX=0withsymmetric windows — correct results with the hierarchical path not selected.
NCCL_GIN_ENABLE, confirming the symmetric-memoryrequirement empirically.
RCCL_HIERARCHICAL_ALLGATHERfor AllGather, and CE vsdefault for AllToAll, at 4/8/16/32 nodes on a release build.
Also tested MI350 and MI300 4 and 8 nodes to compare develop vs develop + my branch to ensure there's no perf regression, sweeps came back clean.
Each arm is gated on the named gtest reporting
[ OK ]and not SKIPPED, datacorrectness, the
[Hierarchical CE]marker, RMA proxy assignment on everyrank, and absence of timeouts or HIP errors.
Test Result
Correctness passes at every scale from 1 to 32 nodes (256 GPUs), with no
skipped tests. Representative 32-node CE AllGather:
Out of bounds values : 0 OK, hierarchical selection markers present, RMAproxy assigned on all 256 ranks, 2560 RMA completions, no warnings.
See ticket for performance impacts.
Submission Checklist