Repository navigation
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
There was a problem hiding this comment.
ARGUS
This PR changes 2,182 lines across 15 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: Worth settling before approval: in the low-CTA SDMA tier only one CTA per rank waits for the inbound puts, so sibling CTAs reduce scratch that has not landed (reduce_scatter.cu:302), and the LSA peer index at reduce_scatter.cu:410 is a world rank.
Comment @argus review to re-run, or @argus re-review after pushing.
| } | ||
| __syncthreads(); | ||
|
|
||
| if (blockIdx.x == (devComm.rank % peerStride)) |
There was a problem hiding this comment.
Only the CTA with blockIdx.x == rank % peerStride waits for the inbound puts, but ncclLsaBarrierSession at line 307 is a cross-rank rendezvous at one barrier index (nccl_device/impl/lsa_barrier__funcs.h:75-79,121-123), not a join across sibling CTAs on this rank. The other CTAs go straight into the reduce at line 317 and read scratchBase with nothing ordering those reads after this rank's puts, and gin.flush is source-side only (nccl_device/gin.h:310-313). Could every CTA wait the signal, or a grid-wide sync go here?
There was a problem hiding this comment.
This was the right call on the most serious item, and it no longer applies to
the PR: the low-CTA SDMA tier came from c588410b5e4, which I dropped. Branch head is now
61a394d42e2, where grep -cE 'sdmaScratchFits|LowCtaSdma|scratchCap' reduce_scatter.cu returns 0.
The shipped kernel is single-tier (direct LSA read-reduce): every CTA joins the one entry
ncclLsaBarrierSession and reads peers' send buffers directly, so there is no staged scratch
whose landing needs ordering.
Worth recording that your analysis matched the hardware. A sweep of that tier did reproduce
in-place data corruption, which is consistent with the ordering hole you describe — sibling
CTAs reducing scratch that hadn't landed, with gin.flush being source-side only. When the
tier returns in its own PR it will carry either a grid-wide sync or a per-CTA signal wait, and
get reviewed on its own.
There was a problem hiding this comment.
Correction to my reply above — on the evidence I cited, not on your finding.
Your reading of c588410b5e4 is confirmed correct by inspection. At that SHA only one block waits:
if (blockIdx.x == (devComm.rank % peerStride))
gin.waitSignal(ncclCoopCta(), signalIndex, signalValue + (uint64_t)nRanks);
gin.flush(ncclCoopCta());
and every block then reads scratchBase in a grid-stride loop (tid = threadIdx.x + blockIdx.x * blockDim.x) after an ncclLsaBarrierSession indexed by blockIdx.x — which, as you said, is a cross-rank rendezvous at one barrier index, not a join across sibling CTAs on this rank. So the hole is real in the code you reviewed.
What I overstated was the hardware corroboration. I wrote that a sweep "did reproduce in-place data corruption, which is consistent with the ordering hole you describe." That sweep ran a different implementation — the hybrid tier on our wip branch, where each block waits for its own nRanks arrivals rather than one block waiting for all — so it is not evidence for the single-waiter hole at this SHA.
A follow-up investigation also re-tallied that corpus and weakened it further: exactly 1 nonzero #wrong row out of 383 datacheck rows (~0.26%), and 105 targeted reruns at the implicated CTA counts reproduced nothing. The earlier ~5% figure I had been working from was the deadlock rate, carried over by mistake. Root cause there is undetermined, and because each row runs the two placements as separate invocations, the in-place specificity rests on n=1 and cannot be called in-place-only.
So your finding stands on its own code reading, which is the part that matters here, and the tier is out of this PR either way. I shouldn't have attached unrelated hardware data to it.
There was a problem hiding this comment.
Resolved by 61a394d42e.
There was a problem hiding this comment.
Agreed on the analysis, and the hazard is real — but the code it describes is the low-CTA SDMA tier that this PR does not ship. The shipped GinReduceScatterKernel is single-tier LSA read-reduce with a per-CTA barrier session (reduce_scatter.cu:271), so no CTA reads scratch another CTA is still filling. Leaving this open deliberately: AICOMRCCL-1803 scopes the -D4 hybrid LSA+GIN variant, and this constraint has to be honoured when that tier lands.
c588410 to
08d1413
Compare
|
@argus re-review |
Addresses the argus review on ROCm#11556. reduce_scatter.cu registered barrier/lsaBarrier/signal pools of max(-V, reduceScatterMaxCtas()) but computed the launched grid with reduceScatterCtas(), which caps an NCCL_GIN_ANVIL_RS_CTAS pin at 128. With -V 8 and RS_CTAS=64 that launched 64 CTAs against 48 pool slots, so blockIdx.x 48..63 indexed past devComm.lsaBarrier. Mirror the AllGather sibling: reduceScatterPoolCtas() is now both the pool size and the ceiling reduceScatterGridCtas() clamps the grid to, so "grid <= pool" is decided in one place. Four new policy unit tests lock the invariant at every band edge (23/23 pass). ReduceScatterParseCtasEnv now rejects a leading '-' and trailing garbage like the AllGather/Broadcast parsers, so RS_CTAS=-2 no longer wraps to a huge unsigned and "8foo" no longer silently pins 8. The two ReduceScatter unit-test targets move out of the AllGather EXISTS guard into their own, keyed on gin_sdma_reducescatter_policy.h; the shared --offload-arch list is hoisted out so both guards can use it. Also adds Broadcast to the device-timing hook comment and the copyright year. Co-authored-by: Cursor <cursoragent@cursor.com>
|
@argus re-review Pushed
Each of your 28 inline comments has a threaded reply stating which of those applies and, where I deferred, why. Two are worth your attention on re-read: the pool-overrun fix added |
There was a problem hiding this comment.
ARGUS
This PR changes 2,126 lines across 16 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: Worth settling before approval: test_ReduceScatterGinSdma.py treats -b as a per-rank size when it is the total, so the CTA-ladder rows and the 2 GiB hang guard run at 1/8th their documented sizes.
Comment @argus review to re-run, or @argus re-review after pushing.
Move MPI launcher assembly, connectivity-gate retry, and Broadcast debug-log capture into gin_sdma_harness.py so AllGather, AllToAll, and Broadcast GIN-SDMA pytest harnesses share one implementation and stop drifting. Co-authored-by: Cursor <cursoragent@cursor.com>
Build device-api CI with ENABLE_DEVICE_API + ENABLE_ROCSHMEM_GIN via cmake, dedupe rocshmem RDC link blocks, move GPU AllGather gtest to gpu-only label, add AllGather policy header missing STATUS, and link rocSHMEM into FixturesDebug for GIN unit tests. Co-authored-by: Cursor <cursoragent@cursor.com>
ENABLE_ROCSHMEM_GIN was defaulting NCCL_GIN_ANVIL_SDMA_ENABLE to 1, so non-GIN fat objects pulled rocSHMEM anvil_device.hpp and failed under RCCL_POISON_HIP_ATOMICS. Keep SDMA opt-in on GIN kernel TUs and rccl-tests RDC binaries. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Device-api CI builds one --rocshmem-gin librccl.so, which leaves GIN rocSHMEM host symbols to the executable. Every binary the matrix launches now whole-archive RDC-links rocSHMEM, not only the GIN collectives. all_reduce_perf still does not build under ENABLE_ROCSHMEM_GIN (same as gin.sbatch), so it is not RDC-linked, not required after `make -k`, and the LSA -D 1/-D 2 suites that launched it are dropped until it builds. The runner's perf-dir sentinel moves to alltoall_perf. Harness: extract gin_perf_argv()/gin_hang_msg() so the three collectives share one argv vector and the AllToAll -A 1 gap is explicit; make detect_ngpus() fail when rocminfo is missing instead of reporting zero GPUs; pin CONN_GATE_RE with a non-gate failure case. Co-authored-by: Cursor <cursoragent@cursor.com>
Parse rocSHMEM's external anvil_device.hpp before RCCL poisons the __hip_atomic_* builtins, while retaining the poison for all following RCCL code. Share the GIN-SDMA backend definitions between regular and DeviceLinker compilation so both pipelines opt in consistently. Restore NCCL_GIN_ANVIL_SDMA_ENABLE=1 on stub-backed tests and let Debug fixtures link rocSHMEM without requiring GIN_ANVIL_UNIT_TESTS. Co-authored-by: Cursor <cursoragent@cursor.com>
Compile a guarded stand-in for rocSHMEM's anvil_device.hpp with a __hip_atomic_load call and verify that the explicitly opted-in external header is parsed before RCCL activates its atomics poison. Co-authored-by: Cursor <cursoragent@cursor.com>
After parsing the external Anvil stub, spell __hip_atomic_load directly in the RCCL probe and require the compiler to reject it. This ensures the external header exception does not disable the poison for following RCCL code. Co-authored-by: Cursor <cursoragent@cursor.com>
…age3e wip Port the single-tier LSA read-reduce kernel, policy/reduce headers, host and GPU unit tests, opt-in pytest regression/devtime smoke, and reduce_scatter_perf rocSHMEM link fix from users/dondai/gin-stage3e-sdma-rs-nccl-2.30.7-wip, reconciled to stage3e conventions (shared CLI devtime, ncclCommQueryProperties devComm reqs, whole-archive rocSHMEM link). Wire ReduceScatter devtime into device-api CI alongside AllToAll. Co-authored-by: Cursor <cursoragent@cursor.com>
…nventions Add initCommConfig and update getDevCommRequirements version guards after cherry-pick onto develop. Co-authored-by: Cursor <cursoragent@cursor.com>
Dispatch pytest matrix entries by collective so RS uses reduce_scatter_perf and RCCL_TESTS_GIN_SDMA_RS, matching the Broadcast GinSdma job. Co-authored-by: Cursor <cursoragent@cursor.com>
Four correctness fixes from the amd-argus review of this PR.
1. reduce_scatter.cu: gate -D 3 on world == LSA team.
The read-reduce calls ncclGetLsaPointer(sendwin, sendoffset, s) with s over
[0, devComm.nRanks), feeding a WORLD rank where an LSA-team index is expected
(ncclGetLsaPointer resolves lsaFlatBase + peer*stride4G). Those coincide only
when the world team is exactly the LSA team, so a comm whose LSA team is a
subset read the wrong peer or outside the aperture. Benign on a single-node
all-P2P comm, which is why the datacheck never caught it. Added the same
nRanks != ncclTeamLsa(comm).nRanks gate all_gather.cu already carries.
2. reduce_scatter.cu: restore the unconditional fp8 prod/mulsum skip.
It had been moved inside `if (deviceImpl != 0)`, which re-enabled fp8 prod and
mulsum on the DEFAULT -D 0 host path that upstream excluded on purpose
(upstream/develop reduce_scatter.cu:121-125). The exclusion is unconditional
again; fp8 avg stays supported, per that upstream comment -- note this is
deliberately narrower than all_reduce.cu, which also excludes fp8 avg. The
device path keeps its separate mulsum-for-any-type skip, which is what it
actually needs since SPECIALIZE_REDUCE_KERNEL returns nullptr there.
3. run-device-api-ci.sh: stop aborting before the failure summary.
run_devtime_smoke ends in `return ${failed}` and was called bare under
`set -euo pipefail`, so any smoke failure exited the script at that call and
the FAILED RUNS summary (and its exit 1) never ran. Call it with `|| true`;
FAILED_RUNS already carries the failures and the summary owns the exit code.
4. run-device-api-ci.sh / gin.sbatch: two config mismatches.
RCCL_TESTS_RS_GIN_TYPE defaulted to 6 (Anvil-SDMA), but device-api.sbatch
builds librccl with a plain `install.sh -lj` so the plugin is compiled out of
the #ifdef ENABLE_ROCSHMEM_GIN entirely. Default to 2 (proxy), matching the
AllToAll smoke in the same script. And gin-tests.json runs
rccl-UnitTestsGinSdmaReduceScatterPolicy while gin.sbatch built only
Fixtures + GinSdmaBroadcastPolicy, so that entry was a permanent SKIP scored
as a failure; it is now built and existence-checked like its sibling.
Not addressed here (style/coverage, not correctness): the repeated three-line
fold, the rsum/rprod widen-narrow duplication, the tautological
gin_sdma_reducescatter_gpu_test.cpp, and the devtime pytest duplication.
GIN_PYTEST_RS_HW_CASES=11 is correct for this branch: the pytest collects 11
cases (8 CTA ladder + 3 hang guard). The review counted 15 against the dropped
low-CTA commit, which added 4 more launches.
Validated: bash -n on all three scripts; policy header compiles with
GIN_SDMA_HOST_ONLY; RS pytest collects 11 tests.
Co-authored-by: Cursor <cursoragent@cursor.com>
Addresses the argus review on ROCm#11556. reduce_scatter.cu registered barrier/lsaBarrier/signal pools of max(-V, reduceScatterMaxCtas()) but computed the launched grid with reduceScatterCtas(), which caps an NCCL_GIN_ANVIL_RS_CTAS pin at 128. With -V 8 and RS_CTAS=64 that launched 64 CTAs against 48 pool slots, so blockIdx.x 48..63 indexed past devComm.lsaBarrier. Mirror the AllGather sibling: reduceScatterPoolCtas() is now both the pool size and the ceiling reduceScatterGridCtas() clamps the grid to, so "grid <= pool" is decided in one place. Four new policy unit tests lock the invariant at every band edge (23/23 pass). ReduceScatterParseCtasEnv now rejects a leading '-' and trailing garbage like the AllGather/Broadcast parsers, so RS_CTAS=-2 no longer wraps to a huge unsigned and "8foo" no longer silently pins 8. The two ReduceScatter unit-test targets move out of the AllGather EXISTS guard into their own, keyed on gin_sdma_reducescatter_policy.h; the shared --offload-arch list is hoisted out so both guards can use it. Also adds Broadcast to the device-timing hook comment and the copyright year. Co-authored-by: Cursor <cursoragent@cursor.com>
parseReduceScatterCtasEnvString and parseReduceScatterUnrollMinBytesString move out of reduce_scatter.cu into gin_sdma_reducescatter_policy.h, behind the host-only guard, so the rejection cases can be locked by the policy unit test without compiling the HIP TU. reduce_scatter.cu keeps the caching wrappers. UNROLL_MIN also gains an overflow guard: a MiB value above (size_t)-1 >> 20 would wrap when shifted, so it now falls back to 0 (tier off) rather than landing on an arbitrary threshold. PoolCoversTheSelfSelectedLadderAtEveryBandEdge now asserts on reduceScatterCtas rather than reduceScatterGridCtas. The clamped form could not exceed its own bound, so the old assertion held by construction; the ladder form tests the property the name describes. reduceScatterGridCtas stays covered by the adjacent pin/clamp tests. Also retires the stale reducescatter-gin-sdma-phase2.md references in favour of AICOMRCCL-1803 for the follow-up put-partials tier. Co-authored-by: Cursor <cursoragent@cursor.com>
ReduceScatterGetCollByteCount divides the -b count by nranks to get the per-rank output slice, so the pytest helpers were understating every size by a factor of nranks. At NP=8 the "256 MiB/rank" hang guard was really 32 MiB/rank and the devtime smoke at 4 MiB never reached the mid CTA band it was meant to exercise. The helpers now take the per-rank slice and multiply by RS_NP, and the devtime smoke moves to 32 MiB so it lands inside [8 MiB, 48 MiB). The datacheck regex was matching the literal "#wrong" column header via the bare "Wrong" alternative, so a clean run could be read as a failure. It now requires "#wrong = <nonzero>". Adds avg and prod to the CTA-ladder matrix, skips the 2 GiB-total guard below RS_NP=8 since that case exists to exercise the 8-way grid-stride path, and raises GIN_PYTEST_RS_HW_CASES from 11 to 15 to match the new launch count. Co-authored-by: Cursor <cursoragent@cursor.com>
61a394d to
6031206
Compare
_assert_datacheck_clean returned early on a clean row and failed on a nonzero one, but output matching none of its patterns fell off the end and passed. A run that died before printing a table row therefore passed on no evidence. The nonzero check now runs first so a nonzero row wins over any clean row printed beside it, and the absence of all three patterns is an explicit failure. Also adds the missing 2026 copyright year to the two ReduceScatter gtest files, matching gin_sdma_reducescatter_policy.h. Co-authored-by: Cursor <cursoragent@cursor.com>
…mments The no-exit-barrier rationale claimed sendbuff is never written by the kernel, which does not hold in place: ReduceScatterGetCollByteCount sets recvInplaceOffset to the slice base, so recvbuff aliases into the sendbuff allocation. The safety argument is that rank R writes only slice R and no peer reads it, so state that instead. SPECIALIZE_REDUCE_KERNEL is defined in reduce_scatter.cu, not common.h. Co-authored-by: Cursor <cursoragent@cursor.com>
dbc66d2 moved the CTA/unroll env parsers into the policy header behind `#if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__)`. HIP parses the whole translation unit in the device pass, including the bodies of host functions it will not codegen, so the guard made name lookup fail inside ReduceScatterParseCtasEnv and ReduceScatterUnrollMinBytes: error: no member named 'parseReduceScatterCtasEnv' in namespace 'gin_sdma_reducescatter' 2 errors generated when compiling for gfx942. That broke reduce_scatter_perf on every arch, not just gfx942. Plain `inline` is already host-only, which is what gin_sdma_allgather_policy.h's parseAllGatherCtasEnv relies on, so the guard is dropped and a comment records why it must not come back. Verified with hipcc --offload-arch=gfx942 -fsyntax-only: the guarded form fails with the errors above, the unguarded form compiles clean. Co-authored-by: Cursor <cursoragent@cursor.com>
There was a problem hiding this comment.
ARGUS
Findings
1 finding(s) that could not be posted as inline comments
projects/rccl/tools/ci/lib/device-api-tests.json:41(Medium) The gating device-api matrix gained a -D 3 AllGather suite at type 6 but no ReduceScatter counterpart The matrix gainsgin-d3-allgatheratNCCL_GIN_TYPE=6but no ReduceScatter counterpart, so the gating device-api job never runs this PR's kernel on Anvil-SDMA, even thoughreduce_scatter_perfis a hard requirement of that job (device-api.sbatch:160-170 exits 1 without it). The only type-6 RS sweep isrccl-gin-rocshmem-sdma-rs-d3in gin-tests.json:66-71, and gin.sbatch:249 turns that job's failures into a non-gating sentinel. Could the RS suite join the matrix next to the AllGather one?
This PR changes 3,730 lines across 29 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: Worth settling before approval: the entry LSA barrier at reduce_scatter.cu:272 syncs relaxed, so the peer read-reduce that follows has no acquire edge.
Comment @argus review to re-run, or @argus re-review after pushing.
| memset(reqs, 0, sizeof(*reqs)); | ||
| switch(deviceImpl) { | ||
| #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) | ||
| case 3: { // single-tier LSA read-reduce: barriers only, no scratch |
There was a problem hiding this comment.
This case 3 returns true unconditionally, so the lsaSize == nRanks gate added at reduce_scatter.cu:167 and the ginType == NCCL_GIN_TYPE_NONE check at :175 both apply only to the >= 2.29.0 overload. ginReduceScatterBody is the same either way and indexes ncclGetLsaPointer(sendwin, sendoffset, s) over [0, devComm.nRanks), which is unchecked arithmetic (lsaFlatBase + peer*stride4G, nccl_device/impl/core__funcs.h:114), so lsaSize < nRanks reads (s - lsaSize) * 4 GiB past the aperture. all_gather.cu's ungated 2.28 arm is safe for a different reason: AllGatherLsaDirect takes lsaRanks and loops lp < lsaRanks (all_gather.cu:163-172, called with lsa.nRanks at :195), so it is bounded in the kernel. Gate: only builds against RCCL 2.28.7 to 2.28.x, and the in-tree RCCL is 2.30.7 (projects/rccl/makefiles/version.mk), so nothing here compiles this arm.
|
|
||
| ncclTeam lsa = ncclTeamLsa(devComm); | ||
| ncclLsaBarrierSession<ncclCoopCta> lsaBar { ncclCoopCta(), devComm, lsa, devComm.lsaBarrier, blockIdx.x }; | ||
| lsaBar.sync(ncclCoopCta(), cuda::memory_order_relaxed); |
There was a problem hiding this comment.
lsaBar.sync(..., cuda::memory_order_relaxed) is the only ordering construct before the peer loads begin at reduce_scatter.cu:316, and it supplies no acquire edge: arrive fences with releaseOrderOf(relaxed) and stores the inbox relaxed, waitInternal loads it with acquireOrderOf(relaxed) (nccl_device/impl/lsa_barrier__funcs.h:72,78,126), and sync is arrive plus wait (:166-169). The comment at :508-510 leans on this barrier to guarantee every peer's sendbuff is filled before any read. The closest sibling is allReduceLsaKernel, which folds peers' sendbuffs the same way and enters with acquire (all_reduce.cu:167 before its fold at :176-179); alltoall.cu:353 does too. The relaxed entries in all_gather.cu:194 and broadcast.cu write to peers and carry the memory_order_release exit sync this kernel deliberately omits.
| lsaBar.sync(ncclCoopCta(), cuda::memory_order_relaxed); | |
| lsaBar.sync(ncclCoopCta(), cuda::memory_order_acquire); |
| RCCL_TESTS_GIN_SDMA_DEVTIME=1 \ | ||
| RCCL_TESTS_RS_EXE="${rs_bin}" \ | ||
| RCCL_TESTS_RS_NP="${NP}" \ | ||
| RCCL_TESTS_RS_GIN_TYPE="${RCCL_TESTS_RS_GIN_TYPE:-2}" \ |
There was a problem hiding this comment.
RCCL_TESTS_RS_GIN_TYPE:-2 binds the proxy backend, while the comment at run-device-api-ci.sh:225-227 says type 6 is in this build and that the default matches gin-tests.json. Three other places say 6: gin-tests.json:69, test_ReduceScatterDevtime.py:41 and its actual default at :83. We asked for 2 last round on the premise that device-api.sbatch built librccl without rocSHMEM; this PR changes that at device-api.sbatch:126 (--rocshmem-gin), so the premise no longer holds and the one gating -D 3 ReduceScatter run never touches Anvil-SDMA.
| RCCL_TESTS_RS_GIN_TYPE="${RCCL_TESTS_RS_GIN_TYPE:-2}" \ | |
| RCCL_TESTS_RS_GIN_TYPE="${RCCL_TESTS_RS_GIN_TYPE:-6}" \ |
| loop = devtimeLoopMid; | ||
| if (devtimeSkipMid >= 0) skip = devtimeSkipMid; | ||
| else skip = (skip < 2) ? skip : 2; | ||
| } |
There was a problem hiding this comment.
This block reproduces gin_devtime::resolveLoopSkip (gin_sdma_devtime.h:34-46, already included at reduce_scatter.cu:20) except that it compares totalBytesCta, the message size summed across ranks, against the 8 MiB and 64 MiB tier constants. AllGatherDeviceTime calls the shared helper with per-rank chunkBytes (all_gather.cu:329-332), and the flag help text says "loop at per-peer >= 8 MiB" (common.cu:2315), so at NP=8 --devtime_loop_mid arms at one eighth of the documented threshold here. Both tier vars default to 0 so nothing arms today.
| } | |
| int loop = 0, skip = 0; | |
| gin_devtime::resolveLoopSkip(count * wordSize(type), loop, skip); |
| *busBw = baseBw * factor; | ||
| switch(deviceImpl) { | ||
| case 3: { // GinReduceScatterKernel: single-tier LSA read-reduce (no scratch) | ||
| if (commProperties.ginType == NCCL_GIN_TYPE_NONE) { |
There was a problem hiding this comment.
This rejects the comm when commProperties.ginType == NCCL_GIN_TYPE_NONE, but ginReduceScatterBody (reduce_scatter.cu:265-515) constructs no ncclGin, issues no put, waits no signal and calls no flush; its only synchronization is the entry ncclLsaBarrierSession at :271. On a P2P-complete node with GIN disabled, reduce_scatter_perf -D 3 therefore refuses a configuration the kernel would run correctly. Is the gate there to hold the door open for the put-partials tier under AICOMRCCL-1803? Please check.
| const int rsNRanks = (rsDc != nullptr) ? rsDc->nRanks : 1; | ||
| const size_t rsTotalBytes = count * (size_t)wordSize(type) * (size_t)rsNRanks; | ||
| static const size_t rsCtasEnv = ReduceScatterParseCtasEnv("NCCL_GIN_ANVIL_RS_CTAS"); | ||
| const int rsGridCtas = gin_sdma_reducescatter::reduceScatterGridCtas( |
There was a problem hiding this comment.
This three-line grid-CTA selection is repeated verbatim at reduce_scatter.cu:639-641, differing only in the total-bytes variable name (rsTotalBytes versus totalBytesCta), and both copies also declare their own function-local static parse of the same env var. The AllGather sibling collapses exactly this into AllGatherResolveLaunchCtas at all_gather.cu:157-160, which both its call sites use. Could a ReduceScatterResolveLaunchCtas(totalBytes) do the same here, leaving one place where pool and grid are reconciled?
| --cmake-options "-DROCSHMEM_INSTALL_DIR=${ROCSHMEM_INSTALL_DIR}" | ||
| ) | ||
| rccl_prefix="${workdir}/projects/rccl/build/release" | ||
| rccl_cfg="$(ls "${rccl_prefix}"/lib*/cmake/rccl/*[Cc]onfig.cmake 2>/dev/null | head -n1)" |
There was a problem hiding this comment.
set -euo pipefail is in force from device-api.sbatch:29, and the exit status of a bare var=$(pipeline) assignment is the command substitution's, which under pipefail is ls's. An unmatched glob makes ls exit 2, so the script terminates at this line and the block at :131-135, the ERROR: no RCCL cmake config message and the ls -laR of what actually exists, never runs. ls stderr is already discarded by 2>/dev/null, so a half-built RCCL ends the job with nothing useful in build-%j.err. Could it take || true?
| rccl_cfg="$(ls "${rccl_prefix}"/lib*/cmake/rccl/*[Cc]onfig.cmake 2>/dev/null | head -n1)" | |
| rccl_cfg="$(ls "${rccl_prefix}"/lib*/cmake/rccl/*[Cc]onfig.cmake 2>/dev/null | head -n1 || true)" |
| _detected_ngpus = detect_ngpus() | ||
| except RuntimeError: | ||
| _detected_ngpus = 0 | ||
| ngpus = max(1, _detected_ngpus) |
There was a problem hiding this comment.
detect_ngpus's docstring says it raises precisely to stop this: a missing binary used to look like zero GPUs because rocminfo | grep | wc -l still exits 0, and callers that clamp with max(1, ...) then silently collect a one-GPU matrix (gin_sdma_harness.py:51-57). Here the raise is caught and the clamp applied anyway, so on a box with no rocminfo and neither visible-devices var ngpus_single collapses to ["1"] where the pre-PR int(math.log2(0)) failed at collection. test_Broadcast.py:45-49 does the same. Could the RuntimeError propagate here instead?
| ngpus = max(1, _detected_ngpus) | |
| ngpus = detect_ngpus() |
| // functions it will not codegen, so hiding these there breaks name lookup in | ||
| // ReduceScatterParseCtasEnv/ReduceScatterUnrollMinBytes and fails the | ||
| // reduce_scatter_perf build on every arch. Plain `inline` is already host-only; | ||
| // it is what gin_sdma_allgather_policy.h's parseAllGatherCtasEnv relies on. |
There was a problem hiding this comment.
This says plain inline is what gin_sdma_allgather_policy.h's parseAllGatherCtasEnv relies on, but that function sits inside #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) at gin_sdma_allgather_policy.h:158, and all_gather.cu:148-155 adds a second device-pass guard at the call site. The reasoning in the two lines above stands on its own; only the appeal to the sibling is false, and a maintainer who trusts it will remove the wrong guard.
| // it is what gin_sdma_allgather_policy.h's parseAllGatherCtasEnv relies on. | |
| // reduce_scatter_perf build on every arch. Plain `inline` is already host-only | |
| // (gin_sdma_allgather_policy.h:158 instead excludes its parsers from that pass). |
| rc, out = _run_devtime(request, device_timing_mode=2) | ||
| assert rc == 0, "reduce_scatter_perf exited {} (mode -B 2)".format(rc) | ||
| _assert_datacheck_clean(out) | ||
| assert "WARN --device_timing=2: no in-kernel device-time" not in out, ( |
There was a problem hiding this comment.
When the timed kernel produces nothing, common.cu:1189-1196 prints this WARN, sets skipMetricRow and the process still exits 0, so assert rc == 0 and _assert_datacheck_clean both pass and only the message wording stands between a silent no-op and a green test. The --devtime_check block that returns testInternalError (common.cu:1216-1220) sits inside the deviceDeltaSec > 0 branch, so the mode-3 test at line 191 has the same hole. Mode 1 at line 175 has the real check, a positive #[rs-devtime] value. Worth dropping these two rather than repairing them?
|
The ReduceScatter slice of this PR is now #12332, rebased onto current Left on their existing PRs: the shared pytest harness (#11894) and the device-api rocSHMEM CI build (#11895). The poison-isolation change stays closed with #11893. The device-api ReduceScatter devtime smoke waits until #11895 lands. |
|
Closing this PR. It conflicts with
The poison-isolation change stays closed with #11893. The device-api ReduceScatter devtime smoke waits until #11895 lands. |
Summary
reduce_scatter_perf -D 3): single-tier LSA read-reduce, host-testable policy, GPU unit tests, and opt-in pytest/devtime smoke.reduce_scatter_perf -D 3, policy gtest,test_ReduceScatterGinSdma.py) and device-api CI devtime smoke.Keep RS single-tier for this PR. A size-hybrid SDMA put-partials path is not required at 4 GiB: the CU LSA pull is already at the xGMI/host-kernel ceiling through 2 GiB (~98–99% of
-D 0), and 4 GiB is the same bandwidth-bound regime (not a new latency/algorithm cliff). Phase-2 pipelined SDMA is a later ~5–9% large-message bet, not a 4 GiB correctness or perf gate.JIRA ID: AICOMRCCL-1803
https://amd-hub.atlassian.net/browse/AICOMRCCL-1803
Test plan
Validated on 8× AMD Instinct MI355X (gfx950), HIP 7.13.99004 / ROCm 7.13.0
Commit validated:
e99ea5c351c.NCCL_GIN_TYPE=6(NCCL_GIN_TYPE_ANVIL_SDMAon thisdevelop-based branch). Plugin load confirmed in-log:
NCCL INFO GIN/Plugin: Assigned plugin gin-anvil-sdma type 6 to comm.All results below exercise the single-tier LSA read-reduce path only — this PR ships
no hybrid/low-CTA tier, and
RS_CTAS/NCCL_GIN_ANVIL_RS_CTASwas never set.projects/rccl-tests/src/reduce_scatter.cucompiles and links with zero warningsor errors;
reduce_scatter_perfbuilt;BUILD SANITY OK(GIN host plugins inlibrccl.so, Anvil-SDMA factory linked).-D 3datacheck, np=8, 128 B → 2 GiB (-b 128 -e 2G -f 2), opsum:float,int32,half— 25/25 size rows each,#wrong= 0 in both theout-of-place and in-place columns,
# Out of bounds values : 0 OK. No hangs; everyrun reached the 2 GiB row and printed its closing summary. Peak 410 GB/s bus bw.
team equals the LSA team (
lsaSize=8==commProperties.nRanks), andGinReduceScatterKernel requires CUDA P2P connectivity across all rankswas neveremitted at np=2, np=4 or np=8. Since the guard runs before the
deviceImplswitch,the successful
-D 3data rows prove it returnedtestSuccess.Not covered: the guard's reject path (world ⊋ LSA team) needs a multi-node comm.
-o all -d allon both paths, notestNotImplementedandno
ncclDevFuncId not foundabort on either:-D 0(default host path): 408 rows, all#wrong= 0. fp8 avg runs; fp8prodand
mulsumskipped;mulsumstill covered for all non-fp8 types.-D 3(GIN device path): 348 rows, all#wrong= 0.mulsumskipped for everytype; fp8
prodskipped; fp8avgruns.-DBUILD_TESTS=ON -DGIN_ANVIL_UNIT_TESTS=ON):rccl-UnitTestsGinSdmaReduceScatterPolicy19/19passed;
rccl-UnitTestsGinSdmaReduceScatterGpu4/4 passed on a real GPU (noskip).
ctest -L unitinrccl-tests: 1/1 passed.Counts are as of
e99ea5c351c. The policy suite has since grown to 25 tests with theReduceScatterPolicyPool(pool/grid clamp) andReduceScatterPolicyEnv(env parserejection) cases; those have not yet been run on hardware — see the MI300X section.
NCCL_GIN_TYPE=5on this branch does not silently pass — it aborts(rc=6, "GIN support is not enabled for this communicator") with zero data rows.
rccl-gin-rocshmem-sdma-rs-d3,rccl-gin-rs-policy-ut,rccl-gin-rs-pytest) — failed at61a394d42e2for a cause outside the RS devicepath, now fixed on this branch. The job never reached any test:
librcclfailed tobuild with
anvil_device.hpp: attempt to use a poisoned identifier(15 errors,gfx950) followed by
ERROR: no librccl.so under .../.ci-out/rccl-install/lib. Rootcause is the upstream pair described under blockers below —
ENABLE_ROCSHMEM_GINdefaults
NCCL_GIN_ANVIL_SDMA_ENABLEto 1 whileRCCL_POISON_HIP_ATOMICSalsodefaults ON.
61a394d42e2carried the poison but not the isolation fix; the branchnow includes
264938bdf36and97377b539e8, so this needs a re-run to confirm.(
device-API / Single node device-api testing, 25m37s, run 34733479064). This jobruns
device-api.sbatch->run-device-api-ci.sh, so it exercises both edits tothat script. From the job log:
test_ReduceScatterDevtime.pyMode1Augment/Mode2DeviceOnly/Mode2WithTimedCheck->3 passed in 67.94s, followed byAll device-api benchmark and devtime smoke runs succeeded.-- that final summaryline is itself the proof the
run_devtime_smokecall no longer aborts the scriptunder
set -euo pipefail.RCCL_TESTS_RS_GIN_TYPEdefaulting to2is alsoconfirmed correct here:
device-api.sbatchbuilds librccl without rocSHMEM, so theAnvil-SDMA plugin is compiled out and only the proxy type can load.
reduce_scatter.cuadditionally compiled clean in this job underENABLE_DEVICE_API=1without rocSHMEM -- a second, independent compile of thecommit in a different configuration from the MI355X build.
Re-validated on 8× AMD Instinct MI300X (gfx942), ROCm 7.0.1 — logs retained
The MI355X figures above predate the barrier-pool clamp and CTA-parser hardening, and that
host has since been repartitioned to
CPX/NPS2, where the Anvil-SDMA context cannot becreated at all (the plugin is assigned, no
context createdline follows, and-D 3correctly refuses with rc=6). So the pool-clamp work was reproduced on MI300X in
SPXinstead. Anvil-SDMA supports MI300X as a first-class path (
anvil.cppcarries an explicitmi300xOamMapfor the MI300X xGMI SDMA layout).Image built from
740d74051e3forgfx942;NCCL_GIN_TYPE=6; device path proven byAssigned plugin gin-anvil-sdma type 6andGIN anvil-sdma: context createdon all8 ranks.
Datacheck,
int32/sum,-D 3 -V 8, out-of-place and in-place — 28/28 rows,#wrong= 0 in every column,Out of bounds values : 0 OKin all four runs.-bis the total send buffer (ReduceScatterGetCollByteCountdivides bynranks), so-b 8M -e 2Gis 1 MiB → 256 MiB per rank:Parser discriminator, 16 MiB/rank (
-b 128M),int32/sum— bandwidthseparates an honoured pin from a rejected one, all with
#wrong= 0:Pre-fix negative control not re-run here. The
RS_CTAS=128hang was observed onMI355X before the clamp; no pre-fix binary was built for gfx942, so that control is
not part of this reproduction.
Policy unit tests not run on MI300X — the image is an
ONLY_FUNCSrccl-testsbuild without
BUILD_TESTS, sorccl-UnitTestsGinSdmaReduceScatterPolicyis absent.Caveats on this reproduction: the image is
gfx942, not thegfx950target; it was builtwith the device linker ON and with rocSHMEM taken from this same monorepo branch (the
Dockerfile default
gin-stage2-gda-sdma-onlydeclaresgin_anvil::sdmawhere this branchexpects
::sdma_anvil); and the policy header carried a local patch equivalent to61b1613ea2c, so the binary is not a byte-clean build of any single pushed commit.Build-configuration blockers hit (pre-existing; not introduced by this PR)
Building
librcclwith-DENABLE_ROCSHMEM_GIN=ONat this branch head needs two flagworkarounds. Both live in RCCL library / rocSHMEM code; this PR touches no file under
projects/rccl/src/and none underprojects/rocshmem/, so neither originates here.-DRCCL_POISON_HIP_ATOMICS=OFF— the default-ONpoison from upstreama171d93c8d1(fix(rccl): replace __hip_atomic builtins with Clang scoped atomics #11321) force-includes
poison_hip_atomics.hinto RCCL's GIN sources, which pull inrocSHMEM headers that still use
__hip_atomic_*(1656 × "attempt to use a poisonedidentifier" in
rocshmem/src/atomic.hppandsrc/sdma/anvil_device.hpp).Upstream's own
projects/rccl/CMakeLists.txtconcedes the conflict — it warns that"ENABLE_ROCSHMEM and ENABLE_ROCSHMEM_GIN are incompatible with RCCL_POISON_HIP_ATOMICS"
— but only warns, because
gin_device_common.hon develop defaultsNCCL_GIN_ANVIL_SDMA_ENABLEto 1 wheneverENABLE_ROCSHMEM_GINis set, pullinganvil_device.hppinto every TU that includesgin__funcs.h.This branch no longer needs the flag workaround:
264938bdf36makes the backendopt-in per TU (default 0) and
97377b539e8parses the external Anvil header ahead of thepoison while keeping it active for all following RCCL code, with
test_poison_hip_atomics.pylocking both directions of that boundary. This is also thecause of the GIN CI failure noted above, which was at a commit predating those two.
install.sh --no-device-linker— thelibrocshmem_gin_host.aextraction(
ld -r+ar,projects/rccl/src/CMakeLists.txt) fails the RDC offload link withlld: error: ...rocshmem_host.bc: Invalid encoding. The surrounding CMake comment onlydocuments the device-linker (default) path, so
ENABLE_DEVICE_LINKER=OFFappearsuntested for
ENABLE_ROCSHMEM_GIN.These were not bisected against pristine
upstream/develop; the attribution is inferredfrom file provenance and this PR's file list. Consequently the validated
librcclis notbit-identical to a default-flag CI build, though nothing in the RS device path depends on
either flag.
Made with Cursor