Repository navigation
fix(devr): size GIN signal shadows for the rounded context count - #13228
Merged
Merged
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
1 task done
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Contributor
Author
|
Follow-ups from the Argus review:
|
1 task done
dlamd1dai
force-pushed
the
users/dondai/rccl-gin-signal-shadow-sizing
branch
from
October 12, 2026 06:41
1e6bf5e to
c5429b5
Compare
Contributor
Author
|
@argus re-review |
Contributor
There was a problem hiding this comment.
ARGUS
Findings
1 finding(s) that could not be posted as inline comments
projects/rccl/test/host/dev-runtime-test.cc:5646(Low) The expected 96 rests on an unstated header default for reqs.ginContextCountGinSignalRegions_UseLargestRoundedContextCountnever setsreqs.ginContextCount, so the whole 96-byte expectation rests on the 4 thatNCCL_DEV_COMM_REQUIREMENTS_INITIALIZERsupplies in nccl_device/host.h:95. That 4 is also what makes the case discriminating, sinceROUNDUP(4, 4)is 4 butROUNDUP(4, 3)is 6. If the public header default is ever retuned this fails as a bare size mismatch with nothing pointing at the cause. Could we pin it in the body?
Bottom line: Looks approval ready.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
Author
|
Pushed
Tested on MI300 (gfx942, ROCm 7.2): @argus re-review |
dlamd1dai
force-pushed
the
users/dondai/rccl-gin-signal-shadow-sizing
branch
from
October 21, 2026 02:26
72cf699 to
f2c7413
Compare
Contributor
Author
|
Rebased onto develop ( @ROCm/rccl-reviewers @ROCm/rccl-leads please review. @argus re-review |
6 tasks done
abouteiller
approved these changes
Oct 23, 2026
abouteiller
left a comment
abouteiller
left a comment
Contributor
There was a problem hiding this comment.
2 minor code sustainability comments. Otherwise lgtm, don't need re-review after done.
dlamd1dai
force-pushed
the
users/dondai/rccl-gin-signal-shadow-sizing
branch
from
October 23, 2026 19:00
f2c7413 to
44388c3
Compare
Contributor
pr-rocm-systems-13228-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-13228-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-13228-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-13228-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-13228-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-13228-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-13228-ruby64-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
The device indexes ginSignalShadows by the devComm's ginContextCount, which ncclGinDevCommSetup rounds up to the chosen backend's connection count, but the shadow region was sized for the requested count. Size it, and the Anvil signal span, for the largest rounded count over every active backend. Co-authored-by: Cursor <cursoragent@cursor.com>
MixedBackends_TakesLargestRoundedCount still passed if the helper kept only the last ROUNDUP, because the furthest-rounding backend was last. Assert the mirrored order so the max accumulator is what the suite locks. Co-authored-by: Cursor <cursoragent@cursor.com>
Drive ncclDevrCommCreateInternal through the resource-window cuMemCreate call and assert that both signal regions use the largest backend-rounded context count. This pins the production integration at dev_runtime.cc:2092-2098, which helper-only tests did not cover. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
…evr sizing Add ncclGinContextCount() in gin_host.h and use it both where ginDevCommSetupWithBackend sets devComm->ginContextCount and where devrGinContextsTotal sizes the per-context resource regions, so the two cannot diverge. Rename nGinContexts to nGinContextsRequested in ncclDevrCommCreateInternal: it only seeds the round-up. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
force-pushed
the
users/dondai/rccl-gin-signal-shadow-sizing
branch
from
October 24, 2026 21:12
44388c3 to
913e54a
Compare
Contributor
pr-rocm-systems-13228-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-13228-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
dlamd1dai
added a commit
that referenced
this pull request
Oct 27, 2026
## Motivation Replaces #13231 with the same commits on an upstream branch, so the RCCL coco PR gate runs without a fork override. The earlier review, including the Argus threads and Aurelien's and Kapil's approvals, is on #13231. Found while reviewing #13228. If `ncclGinConnectOnce` fails, the next `ncclDevrCommCreateInternal` on the same communicator can divide by zero in `ginDevCommSetupWithBackend`. `ncclDevrCommCreateInternal` set `devr->ginEnabled = true` as soon as GIN was requested, before trying to connect, and no failure path cleared it. A connect that fails before `backends[].ginCommCount` is assigned leaves that count at 0, as it was at plugin registration. On retry, `ginEnabled` is already true, so connect is skipped. Setup then runs `ROUNDUP(nContextsTotal, backend->ginCommCount)` with a count of 0. Two argument checks after the latch (a world GIN barrier on a rail connection, and a custom stride of 0) returned with GIN already latched. That left the communicator in the same state. ## Technical Details - `ncclDevrCommCreateInternal` now sets `devr->ginEnabled` only after `ncclGinConnectOnce` succeeds. A failed connect, or an early argument rejection, leaves GIN unlatched, so the next create connects again. `ncclGinConnectOnce` already closes partial connections and leaves `connected` false on failure, so calling it again is safe. - The comment on the latch starts with `RCCL:` and notes that upstream latches before connect, so a future NCCL sync keeps it after connect. - `ginDevCommSetupWithBackend` is unchanged. After this change, a `ginCommCount` of 0 means connect never completed, and clamping it there would hide that. - Test fakes: `dev_runtime_micro_fakes` gets a `g_devrGinConnectOnce` hook, modeled on `g_devrRmaProxyConnectOnce`, so a test can make connect fail. Not changed: if connect succeeds and a later step fails (GIN registration of existing memory, `ncclGinDevCommSetup`, or the window GIN-handle update), GIN stays latched as before. #14231 handles that case. ## Issue Tracking JIRA ID: AICOMRCCL-2879 ## Test Plan - New `DevrCommCreateInternalTest.GinConnectFails_RetryConnectsAgain`: connect fails twice. The test checks that `ginEnabled` stays false and that the second create calls connect again. - New `DevrCommCreateInternalTest.WorldGinBarrierOnRail_RejectedWithoutLatchingGin`: the early rejection does not call connect and does not latch GIN. - Full host-test suite, built host-only with `run_host_tests.sh` in `rocm/dev-ubuntu-22.04:7.2`, after rebasing onto `develop`. ## Test Result - After the rebase: 7909 of 7909 host tests passed across 32 binaries, including both new tests and `develop`'s `GinSignalRegions_UseLargestRoundedContextCount`. - With the `dev_runtime.cc` change reverted, `GinConnectFails_RetryConnectsAgain` fails on the `ginEnabled` check. ## 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> Co-authored-by: Aurelien Bouteiller <aurelien.bouteiller@amd.com> Co-authored-by: amd-argus[bot] <321830313+amd-argus[bot]@users.noreply.github.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
Found while reviewing #12751 (AICOMRCCL-2339). The device indexes
ginSignalShadowsascontextIndex * ginSignalCountfor everycontextIndex < devComm.ginContextCount.ncclGinDevCommSetuproundsginContextCountup to the chosen backend'sginCommCount. InncclDevrCommCreateInternal, though, the shadow region was sized for the requestedreqs->ginContextCount.When the connection count does not divide the requested count, the last contexts' shadows land past the region: in the Anvil signal span, or past the end of the resource window. For example, with
NCCL_GIN_NCONNECTIONS=3and 8 requested contexts, the devComm uses 9 contexts but shadows are allocated for 8.The Anvil span next to it was rounded to the widest active backend, which is also not enough when backends differ:
ROUNDUP(4, 4) = 4 < ROUNDUP(4, 3) = 6.Technical Details
devrGinContextsTotal(ginState, nGinContexts). It returns the maximum ofROUNDUP(nGinContexts, backends[b].ginCommCount)over all active backends, so the size covers whichever backendncclGinDevCommSetuppicks, including a fallback. AginCommCountof 0 is treated as 1.Test Plan
DevrGinContextsTotal.*cases intest/host/dev-runtime-test.cc: no backends, a single backend that rounds up, mixed backends where the widest is not the worst case, and a zero-connection backend.rccl-UnitTestsMicroon MI355 (gfx950, ROCm 7.0.2,ENABLE_ROCSHMEM_GIN=ON).Test Result
rccl-UnitTestsMicro: 1667 of 1667 passed, including the 4 new cases.