Repository navigation
feat(rccl/gin-sdma): GIN SDMA Broadcast (-D3 hybrid Anvil-SDMA) - #11142
Merged
dlamd1dai merged 26 commits intoSep 9, 2026
Merged
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
3 of 5 tasks
![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
Bottom line: Two device-side hazards to settle before merge: the ring constant tables are uploaded only for the first device in a process (broadcast.cu:1214), and every CTA snapshots the GIN signal before a per-block barrier, so a late CTA can wait forever (broadcast.cu:498).
Comment @argus review to re-run, or @argus re-review after pushing.
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 2, 2026
… testable Argus review on ROCm#11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 2, 2026
…sed kernels Argus review on ROCm#11142. The load-bearing fix is the first one; the rest remove the surface that made it hard to see. Per-device ring tables. c_bcastNRings / c_bcastRingSucc / c_bcastRingPos / c_bcastStream are __constant__, so they exist once per DEVICE, but the "already built" gate (g_bcastBuiltN / g_bcastBuiltNRings) was process-wide. In a -g N single-process multi-GPU run device 0 uploaded the decomposition and set the flag, then devices 1..N-1 saw the flag set, skipped their upload, and ran the ring against a zero-filled table: blockIdx.x % c_bcastNRings divides by zero, and every rank reads position 0 and so believes it is the root. Key the state on the current device instead (startColl calls cudaSetDevice immediately before runColl on the device-API path). Distinct host threads own distinct devices, so this also keeps -t N off a shared slot; the remaining shared scratch -- buildBcastRingDecomp's backtracking object and the h_succ/h_pos staging arrays -- becomes thread_local. This was not caught because it needs -g N > 1: the GIN CI entry runs -g 1 under mpirun, and the MI355X gates were MPI-based too. g_bcastLLHandle / g_bcastLLReq become thread_local for the same -t N reason: getDevCommRequirements runs once per thread and links g_bcastLLReq into that thread's requirements list, and one shared node spliced into N lists corrupts them. Note the handle itself did NOT need per-device treatment, contrary to the review comment: ncclDevResourceHandle is a uint32 resource index, and every devComm is created from one identical reqs, so the index is valid on all of them. Ring variants. Five ring kernels shipped but only GinRingSmTableBroadcastKernel was reachable under defaults; the others needed env knobs nothing in the tree sets, and the comments above them already recorded the A/B results that ruled them out. Keep the table kernel and the coprime-stride fallback that N=4 and N=6 genuinely need (no edge-disjoint decomposition exists there), and drop the SDMA-put ring, the single-orientation SM ring, and the point-to-point ring along with NCCL_GIN_ANVIL_BCAST_RING_{SM,MULTI,P2P}. The measurements that justify the choice are kept in the consolidated comment. This also removes the unvalidated ginHandles[0] cast the P2P kernel did without the null/magic checks BroadcastGetSdmaThreshold performs, and collapses the devtime gate's disagreement with the dispatch: builtNRings > 0 is now exactly the condition under which the table kernel runs, so the timed kernel can no longer report a number for a kernel that did not execute. bcastRingCtas(). Used strtol(e, nullptr, 0), so "64foo" parsed as 64 and "010" as octal 8 -- exactly what BroadcastParseCtasEnv exists to reject for the sibling variable. Parse base 10 with a full-string check, and cache the result: it was re-reading getenv at every call site, so the value that sized lsaBarrierCount and the value that picked the launched grid were independent reads of a mutable environment, and a grid larger than the barrier pool is the documented corrupt-or-hang condition. LL cap. The kernel evaluated eligibility against gin_sdma::kBroadcastLLMaxBytes (the 64 KiB compile ceiling) while the host sized the pool and grid from the env-resolved cap (2 KiB by default). They agreed only because bcastLLEligible also requires msgBytes/8 <= nSlots, which re-imposed the cap via the slot count. Pass the resolved cap so both sides evaluate the same bound. Tier selection now goes through the shared helpers: the hybrid kernel switches on bcastKernelTier instead of re-deriving the ladder inline (the header claimed the kernels called it; they did not), the dispatch resolves one bcastLargeTier instead of testing two predicates in an order nothing pinned, and both ring kernels slice with ringStripe/ringChunk. Also: drop the now-unused testLaunchDeviceKernelThreshold launcher and the gin_anvil_sdma.h include that only signalPeer needed; rename the surviving launcher's trailing parameter to kernelArg, since the ring dispatches were passing nChunks into a slot declared sdmaThresholdOverride; remove the dead #else arm on BC_HAVE_LL, whose condition is implied by this block's; and record why the 2.28 requirements arm provisions no LL resource. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 2, 2026
… testable Argus review on ROCm#11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit d64f37c)
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 2, 2026
…sed kernels Argus review on ROCm#11142. The load-bearing fix is the first one; the rest remove the surface that made it hard to see. Per-device ring tables. c_bcastNRings / c_bcastRingSucc / c_bcastRingPos / c_bcastStream are __constant__, so they exist once per DEVICE, but the "already built" gate (g_bcastBuiltN / g_bcastBuiltNRings) was process-wide. In a -g N single-process multi-GPU run device 0 uploaded the decomposition and set the flag, then devices 1..N-1 saw the flag set, skipped their upload, and ran the ring against a zero-filled table: blockIdx.x % c_bcastNRings divides by zero, and every rank reads position 0 and so believes it is the root. Key the state on the current device instead (startColl calls cudaSetDevice immediately before runColl on the device-API path). Distinct host threads own distinct devices, so this also keeps -t N off a shared slot; the remaining shared scratch -- buildBcastRingDecomp's backtracking object and the h_succ/h_pos staging arrays -- becomes thread_local. This was not caught because it needs -g N > 1: the GIN CI entry runs -g 1 under mpirun, and the MI355X gates were MPI-based too. g_bcastLLHandle / g_bcastLLReq become thread_local for the same -t N reason: getDevCommRequirements runs once per thread and links g_bcastLLReq into that thread's requirements list, and one shared node spliced into N lists corrupts them. Note the handle itself did NOT need per-device treatment, contrary to the review comment: ncclDevResourceHandle is a uint32 resource index, and every devComm is created from one identical reqs, so the index is valid on all of them. Ring variants. Five ring kernels shipped but only GinRingSmTableBroadcastKernel was reachable under defaults; the others needed env knobs nothing in the tree sets, and the comments above them already recorded the A/B results that ruled them out. Keep the table kernel and the coprime-stride fallback that N=4 and N=6 genuinely need (no edge-disjoint decomposition exists there), and drop the SDMA-put ring, the single-orientation SM ring, and the point-to-point ring along with NCCL_GIN_ANVIL_BCAST_RING_{SM,MULTI,P2P}. The measurements that justify the choice are kept in the consolidated comment. This also removes the unvalidated ginHandles[0] cast the P2P kernel did without the null/magic checks BroadcastGetSdmaThreshold performs, and collapses the devtime gate's disagreement with the dispatch: builtNRings > 0 is now exactly the condition under which the table kernel runs, so the timed kernel can no longer report a number for a kernel that did not execute. bcastRingCtas(). Used strtol(e, nullptr, 0), so "64foo" parsed as 64 and "010" as octal 8 -- exactly what BroadcastParseCtasEnv exists to reject for the sibling variable. Parse base 10 with a full-string check, and cache the result: it was re-reading getenv at every call site, so the value that sized lsaBarrierCount and the value that picked the launched grid were independent reads of a mutable environment, and a grid larger than the barrier pool is the documented corrupt-or-hang condition. LL cap. The kernel evaluated eligibility against gin_sdma::kBroadcastLLMaxBytes (the 64 KiB compile ceiling) while the host sized the pool and grid from the env-resolved cap (2 KiB by default). They agreed only because bcastLLEligible also requires msgBytes/8 <= nSlots, which re-imposed the cap via the slot count. Pass the resolved cap so both sides evaluate the same bound. Tier selection now goes through the shared helpers: the hybrid kernel switches on bcastKernelTier instead of re-deriving the ladder inline (the header claimed the kernels called it; they did not), the dispatch resolves one bcastLargeTier instead of testing two predicates in an order nothing pinned, and both ring kernels slice with ringStripe/ringChunk. Also: drop the now-unused testLaunchDeviceKernelThreshold launcher and the gin_anvil_sdma.h include that only signalPeer needed; rename the surviving launcher's trailing parameter to kernelArg, since the ring dispatches were passing nChunks into a slot declared sdmaThresholdOverride; remove the dead #else arm on BC_HAVE_LL, whose condition is implied by this block's; and record why the 2.28 requirements arm provisions no LL resource. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 951b01b)
Contributor
Author
Review round: all 29 threads answeredThanks @amd-argus — this was a genuinely useful review. I have replied to every thread individually. Summary of where things stand at head 15 findings are fixed, across two commits:
One finding is a confirmed blocker and is not yet fixed. The 13 remain open and are tracked in their threads: the multi-node LSA-team gate, The three red checks are all externalNone of them are caused by this PR, and none can be cleared from this branch. I have filed both:
Hardware validation at this head8× MI355X, rebuilt from this branch's sources rather than a prebuilt image: full root sweep (roots 0–7) over 128 B – 128 MiB with zero |
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 2, 2026
…to block 0 Both GIN Broadcast kernels sampled their completion signal on every CTA before the entry barrier, while only block 0 ever issues the puts: the put loops walk ranks (r < nRanks / r < N) with a grid-wide tid, so only tid < nRanks executes a body, and nRanks is far below blockDim.x. The entry barrier does not close that gap. ncclBarrierSession is constructed with blockIdx.x, so it orders block n against block n of every peer and nothing else. Block 0 of the root and block n of a non-root share no barrier edge, so the root's put can land before that CTA has run readSignal. It then samples base+1 and waits for base+2, which nobody posts, and the kernel hangs. Moving the read after the barrier does not help, for the same reason. Confining the read, the puts and the wait to block 0 does: block 0's read precedes the block-0-to-block-0 barrier, which precedes the root's put. Kernel completion still covers the wait, because the kernel does not retire until block 0 does. GinScatterAllgatherBroadcastKernel had the identical shape at both of its signals, so the same treatment is applied to the scatter wait and the gather wait, along with the trailing flush -- block 0 issued every put, so it owns the only dirty queues. The put loops now stride by blockDim.x instead of grid-wide, so block 0 still covers every rank if nRanks ever exceeds one block's width. Other CTAs continue to run the grid-strided local copy and must still take the entry barrier, which is a cross-rank barrier keyed by blockIdx.x. Reported-by: amd-argus in ROCm#11142 review Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 2, 2026
…to block 0 Both GIN Broadcast kernels sampled their completion signal on every CTA before the entry barrier, while only block 0 ever issues the puts: the put loops walk ranks (r < nRanks / r < N) with a grid-wide tid, so only tid < nRanks executes a body, and nRanks is far below blockDim.x. The entry barrier does not close that gap. ncclBarrierSession is constructed with blockIdx.x, so it orders block n against block n of every peer and nothing else. Block 0 of the root and block n of a non-root share no barrier edge, so the root's put can land before that CTA has run readSignal. It then samples base+1 and waits for base+2, which nobody posts, and the kernel hangs. Moving the read after the barrier does not help, for the same reason. Confining the read, the puts and the wait to block 0 does: block 0's read precedes the block-0-to-block-0 barrier, which precedes the root's put. Kernel completion still covers the wait, because the kernel does not retire until block 0 does. GinScatterAllgatherBroadcastKernel had the identical shape at both of its signals, so the same treatment is applied to the scatter wait and the gather wait, along with the trailing flush -- block 0 issued every put, so it owns the only dirty queues. The put loops now stride by blockDim.x instead of grid-wide, so block 0 still covers every rank if nRanks ever exceeds one block's width. Other CTAs continue to run the grid-strided local copy and must still take the entry barrier, which is a cross-rank barrier keyed by blockIdx.x. Reported-by: amd-argus in ROCm#11142 review Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 89dd326)
Contributor
Author
Review round 2: test-coverage findings addressed (
|
| run | exit | gin-anvil-sdma markers |
result rows |
|---|---|---|---|
-D 3, GIN Anvil-SDMA |
0 | 16 | 1 |
-D 0, GIN disabled |
0 | 0 | 1 |
A host-path run with GIN switched off exits 0 and prints a full results table. The old assertion could not distinguish it from the path under test.
What changed
- Every measured row must report
#wrong == 0, and must report it as a number rather thanN/A, so a run with data checking disabled is rejected instead of passing silently. At least one measured row must exist, and the out-of-bounds count must be 0. - The run must show the GIN Anvil-SDMA backend actually binding.
_DATA_FAIL_REmatched the literalWrong, which appears in the#wrongcolumn header of every run reaching the results table. It fired on healthy runs, making the connectivity-gate retry beneath it unreachable. It is replaced by_data_failed(), anchored on parsed column values.N/Acounts as unchecked rather than failed, so it surfaces as an assertion instead of consuming the retry budget.- NCCL debug now goes to per-rank files instead of stdout. On stdout the ranks interleave and an INFO line splices into the middle of a results row (
... 60.86 0smci355-...NCCL INFO ...), corrupting the table being parsed. The row regex also accepts the scientific notationgetFloatStrfalls back to, and no longer anchors the line end, where an algo/proto/nchannels group can follow. NCCL_GIN_TYPEwas baked in at 6. TheANVIL_SDMAenum is 6 on develop and a wrong value silently loads no plugin, so it is nowRCCL_TESTS_BCAST_GIN_TYPE(default 6), matching the existingRCCL_TESTS_A2A_GIN_TYPE.- Nothing exported
RCCL_TESTS_GIN_SDMA_BCAST, so the cases skipped wherever they ran. It is exported fromrun-device-api-ci.sh, which is the job that has both GPUs and a pytest runner.lib/gin-tests.jsononly understands therocshmem,rccl-testsandfixtureskinds and cannot drive pytest at all, so the new GIN entry there could never have enabled them.
Validation — all 8 GIN cases pass in 44.6s on 8x MI355X (2.30.7 image), covering 256 MiB and 2 GiB across int32/int64/uint8 plus the 2 GiB and 4 GiB guards. Parsing is separately covered by 11 offline cases including the spliced-row and scientific-notation shapes.
Still open: the multi-node LSA-team gate, moving BcastRingDecomp into the policy header with tests, and the smaller cleanups (baseCount dedupe, sdmaThresholdOverride rename, CMake else() diagnostic, stream-knob documentation).
Not caused by this branch: the two build failures and the therock-pr-bot timeout are tracked as #11155 and #11156; both reproduce on develop independently of this PR.
wenkaidu
reviewed
Sep 3, 2026
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-tests/src/broadcast.cu:3(Low) broadcast.cu copyright still reads 2019-2022 while every GIN sibling reads 2019-2026
Bottom line: The round-1 signal race is fixed and verified. What still blocks is coverage wiring: nothing exports RCCL_TESTS_GIN_SDMA_BCAST and no CI job builds rccl-UnitTestsGinSdmaBroadcastPolicy, so both new suites run in no job.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
There was a problem hiding this comment.
Automated Review Guard for Upfront Scrutiny
Bottom line: The ring's per-step barrier at broadcast.cu:908 asks for release only, so a forwarding rank gets no acquire before reading peer-written data; that should be settled before merge. The new CI pytest entry also selects zero tests.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
Author
Review round 3 addressed (
|
| Finding | Fix |
|---|---|
-k gin_sdma selects nothing (GIN CI exit 5) |
-k GinSdma — pytest case-folds both sides, so gin_sdma ≠ GinSdma |
bcastRingCtas int narrowing |
Clamp in long before (int) cast |
parseSize leading whitespace |
Trim spaces/tabs before sign check; test for " -4K" |
cudaMemcpyToSymbol vs launch stream |
cudaMemcpyToSymbolAsync(..., stream) for all four ring tables |
bcastPoolCtas untested |
BcastPoolCtas.HybridPoolAndRingPeak |
bcastRingChunks only clamps |
Interior case: 64 MiB / 8 CTAs → 128 chunks |
bcastSagCtas env pin untested |
BcastPolicyCtas.SagEnvPinHonoredWithinPool |
pytest import-time int() |
_env_int() falls back on bad values |
| 4 GiB docstring | Says ring path (SAG cannot run at 4 GiB with defaults) |
| Offline parsing guards | Three host-only pytest cases for header / #wrong / sci notation |
Brace style (buildBcastRingDecomp) |
Applied |
Already fixed in earlier commits (replied on threads)
RCCL_TESTS_GIN_SDMA_BCAST/ GIN-job pytest wiring (55e3f180b5e)- Policy UT runs in
Host unit tests (CPU-only)(rccl-UnitTestsGinSdmaBroadcastPolicy, 51 cases) — not the GIN sbatch job, which only buildsrccl-UnitTestsFixtures ENABLE_DEVICE_API=ONingin.sbatch—broadcast_perf -D 3reaches case 3- Multi-node ring gate (
32561e5cb43, @wenkaidu) - Signal race, per-device ring state, pytest assertions (
89dd326a,2df5f510db5)
Deferred (follow-up, not blocking this PR)
- Duplicate ring bodies (
ginRingSmTablevs coprime fallback) — refactor to one shared body; tracked - Shared env parsers (
bcastRingCtasvsBroadcastParseCtasEnv, duplicatedbcastRingMinlambdas) — consolidate with AllGather in AICOMRCCL-2217..2221 stack BcastRingDecomp→ policy header — host-testable, but large move; separate commit- SAG functional pytest case — ring tier is the default large path; SUT sweeps cover SAG; can add
RING_MIN_BYTES=0case later - Ring barrier memory order (
acq_rel) — needs validation against NCCL barrier semantics on MI355X before changing --device_timing+-g N>1ring state — devtime path is opt-in; will key ring upload off launch device in follow-up- Dead
sdmaThresholdOverrideunset path — comment cleanup when threshold helper is unified - Launcher template consolidation — style/readability only
Requesting re-review at head ba186f2fa44.
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 4, 2026
…omment Two Argus review follow-ups on ROCm#11142. No functional change. test_Broadcast.py: _read_debug_logs looped on `path`, shadowing the module-level `path` bound at line 69. 2df5f51 applied the errors="replace" half of the suggestion but left the loop variable, so the finding was still live at head despite being marked resolved. broadcast.cu: the GinHybridBroadcastKernel comment described a fallback to the backend's rsCtx->sdmaThreshold when sdmaThresholdOverride is TEST_SDMA_THRESHOLD_UNSET. That cannot happen from this caller: bcastThr comes from testResolveSdmaThreshold (:1041), and resolveThreshold returns kBroadcastSdmaThresholdDefault rather than the sentinel when both env vars are unset, so BroadcastGetSdmaThreshold is unreachable. The sentinel arm is kept for parity with AllGather; the comment now says that instead of promising a fallback that never runs. Co-authored-by: Cursor <cursoragent@cursor.com>
pvallem
reviewed
Sep 4, 2026
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 4, 2026
_assert_bcast_ok checked that a -D 3 run worked, never that it worked the way the test is named. Every tier produces a correct broadcast, so #wrong stays 0 when a tier stops being reached and its messages fall through to the next one: deleting the ScatterAG arm left test_BroadcastGinSdmaScatterAllgather green with GinScatterAllgatherBroadcastKernel covered by nothing. The two hang guards and the flat segmented tests had the same hole. BroadcastRunColl now names the kernel it launched as "#[bcast-tier] <name>" and the four GIN-SDMA pytest regressions take an expected tier, asserting exactly one ran so a run that switched tiers between roots cannot pass either. The hybrid arm asks gin_sdma::bcastKernelTier -- the policy the kernel itself branches on -- rather than reporting "hybrid", which would leave flat/LSA/LL indistinguishable and the flat tests as uncovered as before. Reported once per distinct tier, since RunColl is called for warmup, every iteration and every root; the dedupe is an atomic fetch_or so -t N host threads cannot race it into printing twice. Opt-in via RCCL_TESTS_BCAST_TIER_LOG, with the env read cached, so default stdout stays byte-identical for the perf harness and the timed path is untouched. The four new offline guards carry GinSdma in their names deliberately: the rccl-gin-bcast-pytest job runs -k GinSdma, and run-gin-ci.sh expands its args unquoted, so a filter containing a space is not available. One of them reproduces the review's scenario directly -- the output a fall-through run would print, which every other assertion accepts. Addresses pvallem's review on ROCm#11142. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
Review round 4 addressed (
|
| tier | reported name |
|---|---|
| pipelined ring, edge-disjoint decomposition | ring-table |
| pipelined ring, no decomposition (N=4, N=6) | ring-multi |
| scatter + allgather | scatter-allgather |
| flat GIN root fan-out | flat-gin |
| direct xGMI LSA stores | lsa-direct |
| LatencyLayover slots | ll |
Three details worth flagging, since they are all places I could have made it weaker:
- It reports the launched kernel, not just
largeTier. You pointed atlargeTierat :1087, which covers ring vs SAG. But the fall-through arm isGinHybridBroadcastKernel, which picks LL / LSA / flat itself, so reporting "hybrid" there would have left the three indistinguishable andtest_BroadcastGinSdmaLargeSegmentedexactly as uncovered as before. So the hybrid arm asksgin_sdma::bcastKernelTier()— the same policy function the kernel branches on — and the ring arm reports table vs multi after the decomposition is built. Every name above maps to one kernel. - Once per distinct tier, not once per launch.
RunCollis called for warmup, every iteration and every root; one line per launch would flood the results table it shares stdout with. The dedupe is an atomicfetch_orso-t Nhost threads can't race it into printing twice. - Opt-in via
RCCL_TESTS_BCAST_TIER_LOG, so default stdout stays byte-identical for the perf harness, and the env read is cached becauseRunCollis on the timed path. Happy to make it unconditional if you'd rather every run be self-describing — see the question at the bottom.
The leading # keeps it clear of the results table (_ROW_RE is anchored on a leading size field), matching the existing #[bcast-devtime] side channel.
Test side — test_Broadcast.py
_assert_bcast_ok(..., expect_tier=) now asserts a tier line exists and that exactly one tier ran and it was an allowed one:
test_BroadcastGinSdmaLargeSegmented→flat-gintest_BroadcastGinSdmaScatterAllgather→scatter-allgathertest_BroadcastGinSdma2GiBHangGuard/4GiB→ring-tableorring-multi
Exactly-one rather than "contains", so a run that took the expected tier for some roots and a different one for the rest doesn't slip through. The ring accepts either name because the decomposition depends on rank count. A missing tier line is a hard failure rather than a skip, so an older binary can't make these pass vacuously.
Four offline guards come with it (no GPU needed), one of which is literally your scenario: it builds the output a fall-through run would produce — correct data, #wrong 0, everything else passing — and asserts the SAG expectation rejects it.
One thing I found on the way
The CI job rccl-gin-bcast-pytest runs -k GinSdma, which never collected the existing offline parsing guards (test_bcast_row_regex_*, test_bcast_data_failed_*) — they run nowhere today. I'd have widened the filter to -k "GinSdma or bcast_", but run-gin-ci.sh:178 expands ${args} unquoted, so the quotes don't survive and pytest gets -k '"GinSdma' plus two stray words. So the new guards are named into the existing filter instead (test_BroadcastGinSdmaTierGuard*) and need no CI change. Glad to rename the four pre-existing ones the same way, or fix the runner's arg handling properly, if you have a preference — didn't want to fold either into this PR uninvited.
pvallem
reviewed
Sep 8, 2026
pvallem
reviewed
Sep 8, 2026
…inSdma The four offline #wrong / scientific-notation cases used snake_case names, so GIN CI's unquoted pytest filter never ran them. Co-authored-by: Cursor <cursoragent@cursor.com>
pvallem
approved these changes
Sep 8, 2026
pvallem
reviewed
Sep 8, 2026
…EOUT GNU timeout killing pytest does not reap the mpirun session; pytest killpg only runs if the inner communicate timeout fires first. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
CI / merge notes at
|
dlamd1dai
added a commit
that referenced
this pull request
Sep 9, 2026
… testable Argus review on #11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
that referenced
this pull request
Sep 9, 2026
…sed kernels Argus review on #11142. The load-bearing fix is the first one; the rest remove the surface that made it hard to see. Per-device ring tables. c_bcastNRings / c_bcastRingSucc / c_bcastRingPos / c_bcastStream are __constant__, so they exist once per DEVICE, but the "already built" gate (g_bcastBuiltN / g_bcastBuiltNRings) was process-wide. In a -g N single-process multi-GPU run device 0 uploaded the decomposition and set the flag, then devices 1..N-1 saw the flag set, skipped their upload, and ran the ring against a zero-filled table: blockIdx.x % c_bcastNRings divides by zero, and every rank reads position 0 and so believes it is the root. Key the state on the current device instead (startColl calls cudaSetDevice immediately before runColl on the device-API path). Distinct host threads own distinct devices, so this also keeps -t N off a shared slot; the remaining shared scratch -- buildBcastRingDecomp's backtracking object and the h_succ/h_pos staging arrays -- becomes thread_local. This was not caught because it needs -g N > 1: the GIN CI entry runs -g 1 under mpirun, and the MI355X gates were MPI-based too. g_bcastLLHandle / g_bcastLLReq become thread_local for the same -t N reason: getDevCommRequirements runs once per thread and links g_bcastLLReq into that thread's requirements list, and one shared node spliced into N lists corrupts them. Note the handle itself did NOT need per-device treatment, contrary to the review comment: ncclDevResourceHandle is a uint32 resource index, and every devComm is created from one identical reqs, so the index is valid on all of them. Ring variants. Five ring kernels shipped but only GinRingSmTableBroadcastKernel was reachable under defaults; the others needed env knobs nothing in the tree sets, and the comments above them already recorded the A/B results that ruled them out. Keep the table kernel and the coprime-stride fallback that N=4 and N=6 genuinely need (no edge-disjoint decomposition exists there), and drop the SDMA-put ring, the single-orientation SM ring, and the point-to-point ring along with NCCL_GIN_ANVIL_BCAST_RING_{SM,MULTI,P2P}. The measurements that justify the choice are kept in the consolidated comment. This also removes the unvalidated ginHandles[0] cast the P2P kernel did without the null/magic checks BroadcastGetSdmaThreshold performs, and collapses the devtime gate's disagreement with the dispatch: builtNRings > 0 is now exactly the condition under which the table kernel runs, so the timed kernel can no longer report a number for a kernel that did not execute. bcastRingCtas(). Used strtol(e, nullptr, 0), so "64foo" parsed as 64 and "010" as octal 8 -- exactly what BroadcastParseCtasEnv exists to reject for the sibling variable. Parse base 10 with a full-string check, and cache the result: it was re-reading getenv at every call site, so the value that sized lsaBarrierCount and the value that picked the launched grid were independent reads of a mutable environment, and a grid larger than the barrier pool is the documented corrupt-or-hang condition. LL cap. The kernel evaluated eligibility against gin_sdma::kBroadcastLLMaxBytes (the 64 KiB compile ceiling) while the host sized the pool and grid from the env-resolved cap (2 KiB by default). They agreed only because bcastLLEligible also requires msgBytes/8 <= nSlots, which re-imposed the cap via the slot count. Pass the resolved cap so both sides evaluate the same bound. Tier selection now goes through the shared helpers: the hybrid kernel switches on bcastKernelTier instead of re-deriving the ladder inline (the header claimed the kernels called it; they did not), the dispatch resolves one bcastLargeTier instead of testing two predicates in an order nothing pinned, and both ring kernels slice with ringStripe/ringChunk. Also: drop the now-unused testLaunchDeviceKernelThreshold launcher and the gin_anvil_sdma.h include that only signalPeer needed; rename the surviving launcher's trailing parameter to kernelArg, since the ring dispatches were passing nChunks into a slot declared sdmaThresholdOverride; remove the dead #else arm on BC_HAVE_LL, whose condition is implied by this block's; and record why the 2.28 requirements arm provisions no LL resource. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
that referenced
this pull request
Sep 9, 2026
…to block 0 Both GIN Broadcast kernels sampled their completion signal on every CTA before the entry barrier, while only block 0 ever issues the puts: the put loops walk ranks (r < nRanks / r < N) with a grid-wide tid, so only tid < nRanks executes a body, and nRanks is far below blockDim.x. The entry barrier does not close that gap. ncclBarrierSession is constructed with blockIdx.x, so it orders block n against block n of every peer and nothing else. Block 0 of the root and block n of a non-root share no barrier edge, so the root's put can land before that CTA has run readSignal. It then samples base+1 and waits for base+2, which nobody posts, and the kernel hangs. Moving the read after the barrier does not help, for the same reason. Confining the read, the puts and the wait to block 0 does: block 0's read precedes the block-0-to-block-0 barrier, which precedes the root's put. Kernel completion still covers the wait, because the kernel does not retire until block 0 does. GinScatterAllgatherBroadcastKernel had the identical shape at both of its signals, so the same treatment is applied to the scatter wait and the gather wait, along with the trailing flush -- block 0 issued every put, so it owns the only dirty queues. The put loops now stride by blockDim.x instead of grid-wide, so block 0 still covers every rank if nRanks ever exceeds one block's width. Other CTAs continue to run the grid-strided local copy and must still take the entry barrier, which is a cross-rank barrier keyed by blockIdx.x. Reported-by: amd-argus in #11142 review Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
that referenced
this pull request
Sep 9, 2026
…omment Two Argus review follow-ups on #11142. No functional change. test_Broadcast.py: _read_debug_logs looped on `path`, shadowing the module-level `path` bound at line 69. 2df5f51 applied the errors="replace" half of the suggestion but left the loop variable, so the finding was still live at head despite being marked resolved. broadcast.cu: the GinHybridBroadcastKernel comment described a fallback to the backend's rsCtx->sdmaThreshold when sdmaThresholdOverride is TEST_SDMA_THRESHOLD_UNSET. That cannot happen from this caller: bcastThr comes from testResolveSdmaThreshold (:1041), and resolveThreshold returns kBroadcastSdmaThresholdDefault rather than the sentinel when both env vars are unset, so BroadcastGetSdmaThreshold is unreachable. The sentinel arm is kept for parity with AllGather; the comment now says that instead of promising a fallback that never runs. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
that referenced
this pull request
Sep 9, 2026
_assert_bcast_ok checked that a -D 3 run worked, never that it worked the way the test is named. Every tier produces a correct broadcast, so #wrong stays 0 when a tier stops being reached and its messages fall through to the next one: deleting the ScatterAG arm left test_BroadcastGinSdmaScatterAllgather green with GinScatterAllgatherBroadcastKernel covered by nothing. The two hang guards and the flat segmented tests had the same hole. BroadcastRunColl now names the kernel it launched as "#[bcast-tier] <name>" and the four GIN-SDMA pytest regressions take an expected tier, asserting exactly one ran so a run that switched tiers between roots cannot pass either. The hybrid arm asks gin_sdma::bcastKernelTier -- the policy the kernel itself branches on -- rather than reporting "hybrid", which would leave flat/LSA/LL indistinguishable and the flat tests as uncovered as before. Reported once per distinct tier, since RunColl is called for warmup, every iteration and every root; the dedupe is an atomic fetch_or so -t N host threads cannot race it into printing twice. Opt-in via RCCL_TESTS_BCAST_TIER_LOG, with the env read cached, so default stdout stays byte-identical for the perf harness and the timed path is untouched. The four new offline guards carry GinSdma in their names deliberately: the rccl-gin-bcast-pytest job runs -k GinSdma, and run-gin-ci.sh expands its args unquoted, so a filter containing a space is not available. One of them reproduces the review's scenario directly -- the output a fall-through run would print, which every other assertion accepts. Addresses pvallem's review on #11142. Co-authored-by: Cursor <cursoragent@cursor.com>
This was referenced Sep 9, 2026
alex-breslow-amd
added a commit
that referenced
this pull request
Sep 10, 2026
This reverts the 26 commits landed by PR #11142 (1d5d2fc..1508be5), restoring the 11 affected files to their state at e9702b4. Reverted commits: 1d5d2fc rccl-tests/gin-sdma: port Broadcast -D 3 from stage3c-sdma-bcast 16da5c0 rccl/gin-sdma: add Broadcast policy unit test and GIN CI entry a847f26 rccl-tests/gin-sdma: align broadcast_perf with NCCL 2.30.7 testEngine API. 47d85d5 rccl-tests/gin-sdma: add 4 GiB Broadcast GIN hybrid hang guard pytest. 94f6e7a style(rccl-tests): apply black formatting to test_Broadcast.py. d6f62cb fix(rccl-tests/gin-sdma): reject negative Broadcast threshold env values f002876 fix(rccl-tests/gin-sdma): align Broadcast with merged AllGather conventions db9ad6d style(rccl-tests/gin-sdma): match merged AllGather testEngine init form d13edf8 fix(rccl-tests/gin-sdma): gate Broadcast RunColl case 3 on ENABLE_DEVICE_API 71c05f9 fix(rccl-tests/gin-sdma): select Anvil-SDMA with NCCL_GIN_TYPE=6 on develop 25df448 style(rccl-tests/gin-sdma): escape the regex backslash in test_Broadcast.py e64f47c rccl-tests/gin-sdma: make the shipped Broadcast tiling and tier order testable 185b45f rccl-tests/gin-sdma: fix Broadcast per-device ring state and trim unused kernels dc19eaf fix(rccl-tests/gin-sdma): confine Broadcast GIN signal read and wait to block 0 9ea1e59 test(rccl-tests/gin-sdma): make Broadcast GIN pytest cases assert and run in CI 72d015a fix(rccl-tests/gin-sdma): gate the Broadcast ring tier on the LSA team 7bba6ea ci(rccl): drop the Broadcast GIN pytest smoke from the device-api job 6f1f937 ci(rccl/gin): run Broadcast GIN pytest from the GIN CI job d8d6678 fix(rccl/gin-sdma): address Argus round-3 review findings 9eaec96 fix(rccl/gin-sdma): SAG pytest, policy UT in GIN CI, ring acq_rel barriers eab9239 fix(rccl/gin-sdma): HIP MemcpyToSymbolAsync needs explicit memcpy kind c81e6d8 fix(rccl/gin-sdma): template bcastMemcpyConstantAsync for scalar symbols 90f2b91 fix(rccl/gin-sdma): unshadow debug log path, correct dead threshold comment 0e06d47 rccl-tests/gin-sdma: assert which Broadcast tier ran b3affe7 test(rccl-tests/gin-sdma): collect Broadcast parser guards under -k GinSdma 1508be5 ci(rccl/gin): keep Broadcast pytest inner budget under GIN_PYTEST_TIMEOUT Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
… testable Argus review on ROCm#11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
…sed kernels Argus review on ROCm#11142. The load-bearing fix is the first one; the rest remove the surface that made it hard to see. Per-device ring tables. c_bcastNRings / c_bcastRingSucc / c_bcastRingPos / c_bcastStream are __constant__, so they exist once per DEVICE, but the "already built" gate (g_bcastBuiltN / g_bcastBuiltNRings) was process-wide. In a -g N single-process multi-GPU run device 0 uploaded the decomposition and set the flag, then devices 1..N-1 saw the flag set, skipped their upload, and ran the ring against a zero-filled table: blockIdx.x % c_bcastNRings divides by zero, and every rank reads position 0 and so believes it is the root. Key the state on the current device instead (startColl calls cudaSetDevice immediately before runColl on the device-API path). Distinct host threads own distinct devices, so this also keeps -t N off a shared slot; the remaining shared scratch -- buildBcastRingDecomp's backtracking object and the h_succ/h_pos staging arrays -- becomes thread_local. This was not caught because it needs -g N > 1: the GIN CI entry runs -g 1 under mpirun, and the MI355X gates were MPI-based too. g_bcastLLHandle / g_bcastLLReq become thread_local for the same -t N reason: getDevCommRequirements runs once per thread and links g_bcastLLReq into that thread's requirements list, and one shared node spliced into N lists corrupts them. Note the handle itself did NOT need per-device treatment, contrary to the review comment: ncclDevResourceHandle is a uint32 resource index, and every devComm is created from one identical reqs, so the index is valid on all of them. Ring variants. Five ring kernels shipped but only GinRingSmTableBroadcastKernel was reachable under defaults; the others needed env knobs nothing in the tree sets, and the comments above them already recorded the A/B results that ruled them out. Keep the table kernel and the coprime-stride fallback that N=4 and N=6 genuinely need (no edge-disjoint decomposition exists there), and drop the SDMA-put ring, the single-orientation SM ring, and the point-to-point ring along with NCCL_GIN_ANVIL_BCAST_RING_{SM,MULTI,P2P}. The measurements that justify the choice are kept in the consolidated comment. This also removes the unvalidated ginHandles[0] cast the P2P kernel did without the null/magic checks BroadcastGetSdmaThreshold performs, and collapses the devtime gate's disagreement with the dispatch: builtNRings > 0 is now exactly the condition under which the table kernel runs, so the timed kernel can no longer report a number for a kernel that did not execute. bcastRingCtas(). Used strtol(e, nullptr, 0), so "64foo" parsed as 64 and "010" as octal 8 -- exactly what BroadcastParseCtasEnv exists to reject for the sibling variable. Parse base 10 with a full-string check, and cache the result: it was re-reading getenv at every call site, so the value that sized lsaBarrierCount and the value that picked the launched grid were independent reads of a mutable environment, and a grid larger than the barrier pool is the documented corrupt-or-hang condition. LL cap. The kernel evaluated eligibility against gin_sdma::kBroadcastLLMaxBytes (the 64 KiB compile ceiling) while the host sized the pool and grid from the env-resolved cap (2 KiB by default). They agreed only because bcastLLEligible also requires msgBytes/8 <= nSlots, which re-imposed the cap via the slot count. Pass the resolved cap so both sides evaluate the same bound. Tier selection now goes through the shared helpers: the hybrid kernel switches on bcastKernelTier instead of re-deriving the ladder inline (the header claimed the kernels called it; they did not), the dispatch resolves one bcastLargeTier instead of testing two predicates in an order nothing pinned, and both ring kernels slice with ringStripe/ringChunk. Also: drop the now-unused testLaunchDeviceKernelThreshold launcher and the gin_anvil_sdma.h include that only signalPeer needed; rename the surviving launcher's trailing parameter to kernelArg, since the ring dispatches were passing nChunks into a slot declared sdmaThresholdOverride; remove the dead #else arm on BC_HAVE_LL, whose condition is implied by this block's; and record why the 2.28 requirements arm provisions no LL resource. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
…to block 0 Both GIN Broadcast kernels sampled their completion signal on every CTA before the entry barrier, while only block 0 ever issues the puts: the put loops walk ranks (r < nRanks / r < N) with a grid-wide tid, so only tid < nRanks executes a body, and nRanks is far below blockDim.x. The entry barrier does not close that gap. ncclBarrierSession is constructed with blockIdx.x, so it orders block n against block n of every peer and nothing else. Block 0 of the root and block n of a non-root share no barrier edge, so the root's put can land before that CTA has run readSignal. It then samples base+1 and waits for base+2, which nobody posts, and the kernel hangs. Moving the read after the barrier does not help, for the same reason. Confining the read, the puts and the wait to block 0 does: block 0's read precedes the block-0-to-block-0 barrier, which precedes the root's put. Kernel completion still covers the wait, because the kernel does not retire until block 0 does. GinScatterAllgatherBroadcastKernel had the identical shape at both of its signals, so the same treatment is applied to the scatter wait and the gather wait, along with the trailing flush -- block 0 issued every put, so it owns the only dirty queues. The put loops now stride by blockDim.x instead of grid-wide, so block 0 still covers every rank if nRanks ever exceeds one block's width. Other CTAs continue to run the grid-strided local copy and must still take the entry barrier, which is a cross-rank barrier keyed by blockIdx.x. Reported-by: amd-argus in ROCm#11142 review Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
…omment Two Argus review follow-ups on ROCm#11142. No functional change. test_Broadcast.py: _read_debug_logs looped on `path`, shadowing the module-level `path` bound at line 69. 2df5f51 applied the errors="replace" half of the suggestion but left the loop variable, so the finding was still live at head despite being marked resolved. broadcast.cu: the GinHybridBroadcastKernel comment described a fallback to the backend's rsCtx->sdmaThreshold when sdmaThresholdOverride is TEST_SDMA_THRESHOLD_UNSET. That cannot happen from this caller: bcastThr comes from testResolveSdmaThreshold (:1041), and resolveThreshold returns kBroadcastSdmaThresholdDefault rather than the sentinel when both env vars are unset, so BroadcastGetSdmaThreshold is unreachable. The sentinel arm is kept for parity with AllGather; the comment now says that instead of promising a fallback that never runs. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
added a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
_assert_bcast_ok checked that a -D 3 run worked, never that it worked the way the test is named. Every tier produces a correct broadcast, so #wrong stays 0 when a tier stops being reached and its messages fall through to the next one: deleting the ScatterAG arm left test_BroadcastGinSdmaScatterAllgather green with GinScatterAllgatherBroadcastKernel covered by nothing. The two hang guards and the flat segmented tests had the same hole. BroadcastRunColl now names the kernel it launched as "#[bcast-tier] <name>" and the four GIN-SDMA pytest regressions take an expected tier, asserting exactly one ran so a run that switched tiers between roots cannot pass either. The hybrid arm asks gin_sdma::bcastKernelTier -- the policy the kernel itself branches on -- rather than reporting "hybrid", which would leave flat/LSA/LL indistinguishable and the flat tests as uncovered as before. Reported once per distinct tier, since RunColl is called for warmup, every iteration and every root; the dedupe is an atomic fetch_or so -t N host threads cannot race it into printing twice. Opt-in via RCCL_TESTS_BCAST_TIER_LOG, with the env read cached, so default stdout stays byte-identical for the perf harness and the timed path is untouched. The four new offline guards carry GinSdma in their names deliberately: the rccl-gin-bcast-pytest job runs -k GinSdma, and run-gin-ci.sh expands its args unquoted, so a filter containing a space is not available. One of them reproduces the review's scenario directly -- the output a fall-through run would print, which every other assertion accepts. Addresses pvallem's review on ROCm#11142. Co-authored-by: Cursor <cursoragent@cursor.com>
anugodavar
pushed a commit
that referenced
this pull request
Sep 17, 2026
… testable Argus review on #11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). Co-authored-by: Cursor <cursoragent@cursor.com>
anugodavar
pushed a commit
that referenced
this pull request
Sep 17, 2026
…sed kernels Argus review on #11142. The load-bearing fix is the first one; the rest remove the surface that made it hard to see. Per-device ring tables. c_bcastNRings / c_bcastRingSucc / c_bcastRingPos / c_bcastStream are __constant__, so they exist once per DEVICE, but the "already built" gate (g_bcastBuiltN / g_bcastBuiltNRings) was process-wide. In a -g N single-process multi-GPU run device 0 uploaded the decomposition and set the flag, then devices 1..N-1 saw the flag set, skipped their upload, and ran the ring against a zero-filled table: blockIdx.x % c_bcastNRings divides by zero, and every rank reads position 0 and so believes it is the root. Key the state on the current device instead (startColl calls cudaSetDevice immediately before runColl on the device-API path). Distinct host threads own distinct devices, so this also keeps -t N off a shared slot; the remaining shared scratch -- buildBcastRingDecomp's backtracking object and the h_succ/h_pos staging arrays -- becomes thread_local. This was not caught because it needs -g N > 1: the GIN CI entry runs -g 1 under mpirun, and the MI355X gates were MPI-based too. g_bcastLLHandle / g_bcastLLReq become thread_local for the same -t N reason: getDevCommRequirements runs once per thread and links g_bcastLLReq into that thread's requirements list, and one shared node spliced into N lists corrupts them. Note the handle itself did NOT need per-device treatment, contrary to the review comment: ncclDevResourceHandle is a uint32 resource index, and every devComm is created from one identical reqs, so the index is valid on all of them. Ring variants. Five ring kernels shipped but only GinRingSmTableBroadcastKernel was reachable under defaults; the others needed env knobs nothing in the tree sets, and the comments above them already recorded the A/B results that ruled them out. Keep the table kernel and the coprime-stride fallback that N=4 and N=6 genuinely need (no edge-disjoint decomposition exists there), and drop the SDMA-put ring, the single-orientation SM ring, and the point-to-point ring along with NCCL_GIN_ANVIL_BCAST_RING_{SM,MULTI,P2P}. The measurements that justify the choice are kept in the consolidated comment. This also removes the unvalidated ginHandles[0] cast the P2P kernel did without the null/magic checks BroadcastGetSdmaThreshold performs, and collapses the devtime gate's disagreement with the dispatch: builtNRings > 0 is now exactly the condition under which the table kernel runs, so the timed kernel can no longer report a number for a kernel that did not execute. bcastRingCtas(). Used strtol(e, nullptr, 0), so "64foo" parsed as 64 and "010" as octal 8 -- exactly what BroadcastParseCtasEnv exists to reject for the sibling variable. Parse base 10 with a full-string check, and cache the result: it was re-reading getenv at every call site, so the value that sized lsaBarrierCount and the value that picked the launched grid were independent reads of a mutable environment, and a grid larger than the barrier pool is the documented corrupt-or-hang condition. LL cap. The kernel evaluated eligibility against gin_sdma::kBroadcastLLMaxBytes (the 64 KiB compile ceiling) while the host sized the pool and grid from the env-resolved cap (2 KiB by default). They agreed only because bcastLLEligible also requires msgBytes/8 <= nSlots, which re-imposed the cap via the slot count. Pass the resolved cap so both sides evaluate the same bound. Tier selection now goes through the shared helpers: the hybrid kernel switches on bcastKernelTier instead of re-deriving the ladder inline (the header claimed the kernels called it; they did not), the dispatch resolves one bcastLargeTier instead of testing two predicates in an order nothing pinned, and both ring kernels slice with ringStripe/ringChunk. Also: drop the now-unused testLaunchDeviceKernelThreshold launcher and the gin_anvil_sdma.h include that only signalPeer needed; rename the surviving launcher's trailing parameter to kernelArg, since the ring dispatches were passing nChunks into a slot declared sdmaThresholdOverride; remove the dead #else arm on BC_HAVE_LL, whose condition is implied by this block's; and record why the 2.28 requirements arm provisions no LL resource. Co-authored-by: Cursor <cursoragent@cursor.com>
anugodavar
pushed a commit
that referenced
this pull request
Sep 17, 2026
…to block 0 Both GIN Broadcast kernels sampled their completion signal on every CTA before the entry barrier, while only block 0 ever issues the puts: the put loops walk ranks (r < nRanks / r < N) with a grid-wide tid, so only tid < nRanks executes a body, and nRanks is far below blockDim.x. The entry barrier does not close that gap. ncclBarrierSession is constructed with blockIdx.x, so it orders block n against block n of every peer and nothing else. Block 0 of the root and block n of a non-root share no barrier edge, so the root's put can land before that CTA has run readSignal. It then samples base+1 and waits for base+2, which nobody posts, and the kernel hangs. Moving the read after the barrier does not help, for the same reason. Confining the read, the puts and the wait to block 0 does: block 0's read precedes the block-0-to-block-0 barrier, which precedes the root's put. Kernel completion still covers the wait, because the kernel does not retire until block 0 does. GinScatterAllgatherBroadcastKernel had the identical shape at both of its signals, so the same treatment is applied to the scatter wait and the gather wait, along with the trailing flush -- block 0 issued every put, so it owns the only dirty queues. The put loops now stride by blockDim.x instead of grid-wide, so block 0 still covers every rank if nRanks ever exceeds one block's width. Other CTAs continue to run the grid-strided local copy and must still take the entry barrier, which is a cross-rank barrier keyed by blockIdx.x. Reported-by: amd-argus in #11142 review Co-authored-by: Cursor <cursoragent@cursor.com>
anugodavar
pushed a commit
that referenced
this pull request
Sep 17, 2026
…omment Two Argus review follow-ups on #11142. No functional change. test_Broadcast.py: _read_debug_logs looped on `path`, shadowing the module-level `path` bound at line 69. 2df5f51 applied the errors="replace" half of the suggestion but left the loop variable, so the finding was still live at head despite being marked resolved. broadcast.cu: the GinHybridBroadcastKernel comment described a fallback to the backend's rsCtx->sdmaThreshold when sdmaThresholdOverride is TEST_SDMA_THRESHOLD_UNSET. That cannot happen from this caller: bcastThr comes from testResolveSdmaThreshold (:1041), and resolveThreshold returns kBroadcastSdmaThresholdDefault rather than the sentinel when both env vars are unset, so BroadcastGetSdmaThreshold is unreachable. The sentinel arm is kept for parity with AllGather; the comment now says that instead of promising a fallback that never runs. Co-authored-by: Cursor <cursoragent@cursor.com>
anugodavar
pushed a commit
that referenced
this pull request
Sep 17, 2026
_assert_bcast_ok checked that a -D 3 run worked, never that it worked the way the test is named. Every tier produces a correct broadcast, so #wrong stays 0 when a tier stops being reached and its messages fall through to the next one: deleting the ScatterAG arm left test_BroadcastGinSdmaScatterAllgather green with GinScatterAllgatherBroadcastKernel covered by nothing. The two hang guards and the flat segmented tests had the same hole. BroadcastRunColl now names the kernel it launched as "#[bcast-tier] <name>" and the four GIN-SDMA pytest regressions take an expected tier, asserting exactly one ran so a run that switched tiers between roots cannot pass either. The hybrid arm asks gin_sdma::bcastKernelTier -- the policy the kernel itself branches on -- rather than reporting "hybrid", which would leave flat/LSA/LL indistinguishable and the flat tests as uncovered as before. Reported once per distinct tier, since RunColl is called for warmup, every iteration and every root; the dedupe is an atomic fetch_or so -t N host threads cannot race it into printing twice. Opt-in via RCCL_TESTS_BCAST_TIER_LOG, with the env read cached, so default stdout stays byte-identical for the perf harness and the timed path is untouched. The four new offline guards carry GinSdma in their names deliberately: the rccl-gin-bcast-pytest job runs -k GinSdma, and run-gin-ci.sh expands its args unquoted, so a filter containing a space is not available. One of them reproduces the review's scenario directly -- the output a fall-through run would print, which every other assertion accepts. Addresses pvallem's review on #11142. Co-authored-by: Cursor <cursoragent@cursor.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
… testable Argus review on #11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). Co-authored-by: Cursor <cursoragent@cursor.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
… testable Argus review on #11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). Co-authored-by: Cursor <cursoragent@cursor.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
… testable Argus review on #11142 pointed out that several of the policy suite's assertions covered code the kernels do not run, while the code they do run was untested. Address that in the header rather than the tests: * Add ringStripe()/ringChunk(). The default large tier is the SM ring, and its CTA-stripe and pipeline-chunk splits were open-coded at each kernel that walks them, appearing in no test. The one tiling the suite did prove exactly (sagChunk) is used only by the scatter+allgather tier. New cases assert both splits partition their range exactly, stay contiguous from zero, stay balanced to within one element, and compose to cover the buffer for an indivisible (prime) element count. * Add bcastLargeTier(). Under the shipped defaults bcastUseRing (>=32 MiB) and bcastUseScatterAllgather (>=2 MiB) are BOTH true for every large message, so which tier runs is fixed by the order the dispatch tests them and by nothing else. Each predicate was asserted in isolation, and "both true" satisfies all of those assertions, so swapping the two blocks would not have failed a test. Returning the resolved tier as an enum lets the 1 MiB / 2 MiB / 32 MiB boundaries and the ring opt-out pin the precedence directly. * Reject overflowing size values in parseSize(). strtoull saturates at ULLONG_MAX instead of failing, and the suffix multiply can then wrap a huge value back down: "17179869184G" is 2^64 bytes, which wrapped to 0, and 0 is meaningful here (it force-enables GIN). That is the same silent mis-tiering shape as the negative-value defect fixed earlier in this branch. Also pin that "0x10" reads as 0 rather than 16, since the parser is base 10 by contract. * Drop alignChunkCount(). It had no caller in broadcast.cu: both the SAG and ring paths slice with sagChunk, which deliberately does an unaligned even split with the remainder folded into the last rank, so the header was shipping an alignment rule the kernels contradict, plus two test cases and eight assertions backing it. * Make the CTA-ladder assertions relational. Eight of them compared bcastHybridCtas' result against the very constant it returns on that branch, so they pinned which branch was taken but not the ladder's shape. Add a case asserting LL < LSA, SDMA < LSA, and every grid within [1, pool]. Also inline the basis for the header's defaults instead of pointing at ddai-artifacts/docs/gin-anvil-sdma-broadcast-design-plan.md, which is not in the repo and left the bare section numbers in broadcast.cu unresolvable. Host policy suite: 48/48 pass (was 38; -2 dead-helper cases, +12 new). 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.
Summary
Adds the GIN Anvil-SDMA hybrid Broadcast (
broadcast_perf -D 3,NCCL_GIN_TYPE=6), completing the SDMA collective set alongside the merged AllToAll (#7826, #10420) and AllGather (#10930).JIRA ID : AICOMRCCL-1799
Size-adaptive tiering (
-D 3), all selected by a single pure policy header:gin.put()fan-out, receiver-sidewaitSignal(+1)busBw = B_root_egress / N)gin_sdma_broadcast_policy.h: pure, host-testable tier policy — threshold + LL-cap env resolution, scatter/ring chunk math, signal-count sizing. Compiles underGIN_SDMA_HOST_ONLYas plain C++ and is called from the__global__kernels, so the host-tested decisions are exactly the ones that run on device.broadcast.cu: the tiered kernels, devtime hook via the sharedgin_sdma_devtime.hscaffold, rocSHMEM RDC link forbroadcast_perf.rccl-UnitTestsGinSdmaBroadcastPolicyhost GTest (51 cases), plus opt-in pytest multi-segment regressions with 2 GiB / 4 GiB hang guards.broadcast_perf -D 3(rccl-gin-rocshmem-sdma-bcast-d3) andpytest test_Broadcast.py -k gin_sdma(rccl-gin-bcast-pytest) viarun-gin-ci.sh, both withNCCL_GIN_TYPE=6.11 commits, 7 files, +2242/−50. Touches zero AllGather files.
Why this lives in
rccl-testsrather thansrc/algorithms/gin/#10386 recently added an in-library GIN SDMA AllToAll under
projects/rccl/src/algorithms/gin/, so it is fair to ask why Broadcast is not there too.These are different layers rather than competing homes. The in-library collective is host-side dispatch through
collectives.cc: the application callsncclAllToAlland RCCL selects a GIN implementation internally. This PR is the device-API path — the collective runs inside the caller's own kernel viancclDevComm, which is what-D 3in rccl-tests exercises and what the GIN device-kernel work item (AICOMRCCL-1796..1800) is chartered to deliver. It has no host-side entry point to dispatch from.This also follows the merged AllGather precedent (#10930) exactly, keeping the two GIN-SDMA device collectives side by side in one tree so the shared policy/devtime scaffolding stays comparable. Promoting the device-API tiers into
src/algorithms/gin/later is a reasonable direction, but it should be one deliberate migration covering AllGather and Broadcast together, not a divergence introduced here.Base
Rebased onto
developat1c1419a9434(conflict-free;gin-tests.jsonandrccl/test/CMakeLists.txtboth picked up upstream's newer entries alongside this PR's). Previously stacked on the AllGather branch, now merged as #10930 — all inherited AllGather commits were dropped, along with three stale carries that would have regresseddevelop: thegin_anvil_sdma_oss7_device.ccconn-check TU that #10930 deliberately removed (unlinked__hip_fatbinin GIN CI, being restored properly in #11092), analloc.hrevert of the side-stream pooling, and pre-review versions of the AllGather sources.Defects found and fixed while preparing this PR
1.
-D 3failed to compile in the default configuration (42eaaeb2927)BroadcastRunCollcase 3 was gated on the NCCL version alone, but everything it calls —testLaunchDeviceKernelThreshold{,Ctas,LL}, theGin*BroadcastKernelsymbols,buildBcastRingDecompand theg_bcastBuilt*ring state — is defined insidedefined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)blocks.ENABLE_DEVICE_APIdefaults to OFF, so a stock rccl-tests configure against NCCL ≥ 2.28.7 compiled the call site with every callee preprocessed out.Now gated on both macros, matching the merged
all_gather.cuandBroadcastDeviceTimefurther down this same file. Verified by compiling both ways:2. Negative threshold env values silently pinned the LSA tier (
02fc3917082)parseSize()fed the raw string tostrtoull(), which accepts a leading-and wraps it nearSIZE_MAX.-1happened to land onkThresholdUnsetand was harmless, but-4Kparsed as a ~16 EiB threshold, so every message compared as "below threshold", took the LSA tier, and the sweep still reported pass — the GIN/SDMA path would simply never run. Rejects the sign before parsing now, matching the AllGather policy header.3. Broadcast CI and pytest selected the wrong GIN backend (
956427cfbf1)Both requested
NCCL_GIN_TYPE=5. That is correct on the 2.30.7 line this work was developed against, but the enum was renumbered ondevelop:Plugin selection is a strict equality match on
props.netDeviceType, so ondevelopa request for5resolves to the rocSHMEM GDA backend rather than Anvil-SDMA. The Broadcast CI entry would have silently exercised the wrong backend, while sitting directly between two AllToAll entries that already say6. Now6everywhere, matching the mergedtest_AllGather.py,test_AllToAll.pyand the AllToAll CI entries.4.
.getDevCommRequirementsleft null on 2.29+ without the device API (a23b7871504)Wired under the same
NCCL >= 2.29 || (ENABLE_DEVICE_API && >= 2.28)guardcommon.hdeclares the slot with. Same commit hardensBroadcastParseCtasEnvagainst a leading-, trailing junk (8fooparsed as8) and a literal sentinel collision, matchinggin_sdma_allgather::parseAllGatherCtasEnv().Test plan
rccl-UnitTestsGinSdmaBroadcastPolicy— 51/51 pass (host-only;Host unit tests (CPU-only)CI green on prior run)broadcast_perf -D 3formal gate, 8× MI355X — PASS, 336 measured rows across all 8 roots,#wrong=0out-of-place and in-place#wrong=0ENABLE_DEVICE_APIon/off compile check (see above)broadcast_perfentry ingin-tests.jsonviarun-gin-ci.sh(rccl-gin-rocshmem-sdma-bcast-d3,NCCL_GIN_TYPE=6; GIN job re-running at55e3f180b5e)pytest test_Broadcast.py -k gin_sdmaend-to-end — 8/8 pass on 8× MI355X SUT (44.6s); wired into GIN CI asrccl-gin-bcast-pytest(55e3f180b5e: exportsRCCL_TESTS_GIN_SDMA_BCAST=1,BCAST_EXE,BCAST_GIN_TYPE=6; not device-API, which runs PROXY/type 2)SUT evidence — 8× MI355X gfx950,
smci355-ccs-aus-m03-17, Anvil-SDMA backendRun against a docker image built from this commit's Broadcast sources. The image tree is the NCCL 2.30.7 validation line, where the Anvil-SDMA backend is
NCCL_GIN_TYPE=5(it is6ondevelop— see defect 3 above); the runs were confirmed to bind the intended backend viaAssigned GIN plugin gin-anvil-sdma to commin the init log, not by trusting the number.Follow-up (not in this PR)
The three collectives now carry duplicated policy scaffolding: two
__host__ __device__attribute macros (GIN_SDMA_AG_HD/GIN_SDMA_HD), two namespaces (gin_sdma_allgather/gin_sdma), threshold precedence and 16-byte chunk alignment implemented twice, size-string parsing three times, and near-identical pytest launch/retry helpers. Worth consolidating once this lands and after the AICOMRCCL-2217..2221 stack, which already dedups the Anvil-SDMA env parsing and the pytest connectivity-gate retry.Separately,
GIN_SDMA_HDis defined unconditionally ingin_anvil_sdma_put_policy.hwhile the Broadcast header guards it with#ifndef. Safe today only because the expansions are identical; it wants an#ifndefon the canonical definition. Deliberately left out of this PR to keep the diff to Broadcast.