[dlamd1dai]

Summary

• Split the ReduceScatter device path out of feat(rccl-tests): GIN SDMA ReduceScatter (-D 3) #11556 onto current develop. The shared pytest harness stays on rccl-tests: share GIN-SDMA pytest harness #11894, the device-api rocSHMEM CI build stays on ci(rccl): build device-api with rocSHMEM GIN #11895, and the poison-isolation change stays closed with rccl: isolate GIN-SDMA TUs from HIP atomics poison (AICOMRCCL-2221) [1/3] #11893.
• Add a GIN Anvil-SDMA ReduceScatter device path (reduce_scatter_perf -D 3): single-tier LSA read-reduce, host policy unit test, GPU addressing test, opt-in pytest, and GIN CI entries (rccl-gin-rocshmem-sdma-rs-d3, rccl-gin-rs-policy-ut, rccl-gin-rs-pytest).
• The device-api ReduceScatter devtime smoke (test_ReduceScatterDevtime.py and the device-api job hook) waits until ci(rccl): build device-api with rocSHMEM GIN #11895 lands.
• On current develop, NCCL_GIN_TYPE_ANVIL_SDMA is 7 (type 6 is rocSHMEM GDA). The new CI entry and the pytest default use 7. Hardware runs recorded on feat(rccl-tests): GIN SDMA ReduceScatter (-D 3) #11556 used 6 because that number was Anvil SDMA on that older base.

Host-path floating-point avg stays skipped when deviceImpl == 0, matching the skip that landed on develop after #11556 branched. The GIN device path still runs avg and still skips mulsum.

JIRA ID: AICOMRCCL-1803
https://amd-hub.atlassian.net/browse/AICOMRCCL-1803

Test plan

• reduce_scatter_perf builds with ENABLE_DEVICE_API and ENABLE_ROCSHMEM_GIN (RDC-linked like broadcast_perf)
• rccl-UnitTestsGinSdmaReduceScatterPolicy passes (ctest -L unit)
• rccl-UnitTestsGinSdmaReduceScatterGpu passes on a GPU
• RCCL_TESTS_GIN_SDMA_RS=1 pytest test_ReduceScatterGinSdma.py with NCCL_GIN_TYPE=7
• GIN CI runs rccl-gin-rocshmem-sdma-rs-d3, rccl-gin-rs-policy-ut, and rccl-gin-rs-pytest
• Host -D 0 still skips floating avg and still runs mulsum

Made with Cursor

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[dlamd1dai]

Draft for the low-CTA GIN SDMA ReduceScatter tier, stacked on this branch so the diff is only that tier: dlamd1dai#17

It should move to ROCm/rocm-systems develop after this PR merges.

[amd-argus]

ARGUS

Findings

1 finding(s) that could not be posted as inline comments

• projects/rccl/tools/ci/lib/gin.sbatch:205 (Low) gin.sbatch never asserts reduce_scatter_perf was built, so a break in the new RDC link is swallowed gin.sbatch builds rccl-tests with make -k || true at line 198 and then verifies only alltoall_perf (line 200) and broadcast_perf (line 205). reduce_scatter_perf, which this PR newly gives a -fgpu-rdc whole-archive rocSHMEM link, is not checked anywhere in the file, so a compile break in exactly the new path is swallowed by -k and surfaces only as a downstream SKIP ... perf binary not found in an explicitly non-gating job. Could a third copy of the lines 200-204 block cover it at build time?

This PR changes 1,902 lines across 14 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 warp-unroll arm at reduce_scatter.cu:388 ships with no test, and the GPU addressing test derives its expected values the same way as the code it checks.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

[dlamd1dai]

Addressed the review findings in be3c3a6829 and 186dfb76b5:

• removed the untested legacy warp-unroll path and its environment toggle
• shared the ReduceScatter slice-offset helper between production and GPU-test code, with independent expected values and explicit policy coverage
• asserted that reduce_scatter_perf was built after the non-gating make -k
• gave ReduceScatter pytest a dedicated 3600-second outer budget
• rejected whitespace-prefixed negative CTA pins
• removed the inaccurate AllGather barrier comparison

Validation completed: shell syntax checks, policy-header compile/run smoke test, GIN config parser tests, and git diff --check.

@argus re-review

[amd-argus]

ARGUS

This PR changes 1,751 lines across 14 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:244 is memory_order_relaxed before the kernel reads every peer's sendbuff, where all_reduce.cu:167 uses acquire for the same read-reduce shape.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

[amd-argus]

High

This lsaBar.sync is cuda::memory_order_relaxed and is the only synchronization in ginReduceScatterBody before it reads every peer's sendbuff through ncclGetLsaPointer (line 272 onward). acquireOrderOf(relaxed) is relaxed (nccl_device/utility.h:351-355), so the wait-side load carries no acquire and those remote reads are not ordered after observing the peer arrived. all_reduce.cu:167 is the same LSA read-reduce shape and uses acquire; all_gather.cu:194 can stay relaxed because its tier only writes peers and publishes with release at :196. releaseOrderOf(acquire) is relaxed, so the arrive side costs nothing.

Suggested change

	lsaBar.sync(ncclCoopCta(), cuda::memory_order_relaxed);
	lsaBar.sync(ncclCoopCta(), cuda::memory_order_acquire);

[dlamd1dai]

Good catch — fixed in 8efddab141, taken exactly as suggested.

The part that convinced me this is a real bug rather than a defensible relaxation is that the function's own exit comment already asserts the guarantee the relaxed barrier does not provide:

The entry barrier already guarantees every peer's sendbuff is filled before any read.

That is the invariant the "NO exit barrier" argument rests on, and with acquireOrderOf(relaxed) returning relaxed at utility.h:351-355 the wait-side load carried no acquire, so the ncclGetLsaPointer reads from line 272 onward were not ordered after observing the peer's arrival. The comment described the intended design; the code did not implement it.

Your precedent comparison holds up on both sides. all_reduce.cu:167 is the same LSA read-reduce shape and acquires. all_gather.cu:194 can stay relaxed because all_gather.cu:169 uses the LSA pointer as a destination — that tier writes peers rather than reading them — and publishes with the release at :196. ReduceScatter is the read case, so it needs the acquire and has no exit barrier to lean on.

Since -D 3 exists for performance, it seemed worth measuring the cost rather than reasoning about it. At 1 MiB total, five repeats each on 8x MI350X: acquire 10.56, 10.56, 10.69, 10.86, 10.67 us against relaxed 11.13, 10.84, 11.00, 10.67, 10.64 us. The acquire arm is marginally faster on those means, which is to say the difference is entirely run-to-run noise — as expected, since releaseOrderOf(acquire) is relaxed so the arrive side is unchanged and the wait side gains one acquire per CTA per collective. A single earlier sample had suggested an 8% cost; that did not survive repetition.

Datacheck is clean at 4 MiB/rank for float and int32 (Out of bounds values : 0 OK), along with the six new narrowing-dtype cases.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

Medium

rccl-UnitTestsGinSdmaReduceScatterGpu is absent from this target list and has no gin-tests.json entry, so the 191 lines of GPU test this PR adds are never compiled, let alone executed: nothing in RCCL CI runs ctest, as projects/rccl/test/CMakeLists.txt:23 states. The sibling policy UT got both halves in this same PR (this line, plus the rccl-gin-rs-policy-ut entry). Could the GPU target join both? The "gpu"-only label is fine as it stands.

[dlamd1dai]

Good catch — fixed in 8efddab141. The GPU target joins both halves, as you suggested.

The gap holds up end to end: add_executable(rccl-UnitTestsGinSdmaReduceScatterGpu ...) sits at test/CMakeLists.txt:714 with its own add_test and LABELS "gpu", but the target appeared in neither gin.sbatch's build step nor gin-tests.json, and since nothing in RCCL CI runs ctest those 191 lines were compiled by no job and executed by none.

gin.sbatch now builds it alongside rccl-UnitTestsGinSdmaReduceScatterPolicy and fails the job with its own error message if the binary is missing, matching the treatment the policy UT already had. gin-tests.json gains an rccl-gin-rs-gpu-ut entry under kind fixtures, with a comment recording why an entry is needed at all despite the ctest registration:

{
  "_comment": "GPU-labelled in ctest, but RCCL CI does not run ctest, so it needs its own entry to execute at all.",
  "name": "rccl-gin-rs-gpu-ut",
  "kind": "fixtures",
  "bin": "rccl-UnitTestsGinSdmaReduceScatterGpu",
  "env": [],
  "args": ""
}

I left the "gpu" ctest label alone, as you said it reads fine — it still does the right thing for a local ctest -L unit.

One caveat on verification: I confirmed the JSON parses, the parser emits the new row, and gin.sbatch is syntactically valid, but I have not run the full GIN sbatch job, so the first real proof that this target builds under that job's configuration will come from CI.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

Medium

This dtype list plus the hang guard's int64/half at line 195 means bfloat16, fp8_e4m3 and fp8_e5m2 never run through -D 3, and the rccl-gin-rocshmem-sdma-rs-d3 entry passes no -d, so it adds only float. Those three types are what gin_sdma_reduce.h's narrow-back-to-T-on-every-pairwise-step overloads exist for (gin_sdma_reduce.h:51 and :56 for rsum alone), and ReduceScatterRunTest does not exclude them on the device path. Could they join the ladder for sum and avg?

[dlamd1dai]

Good catch, and fixed in 8efddab141. You are right that the three types the narrowing overloads exist for were the three the ladder never reached.

I confirmed the device path accepts them — RS_BF16_CASE and RS_FP8_CASE are wired into SPECIALIZE_REDUCE_KERNEL, and the only exclusion there is prod for the two fp8 types. So rather than widen the existing dtype list, which would have dragged all three through prod, min and max as well, I added a separate case limited to the ops that actually narrow:

@_rs_skip
@pytest.mark.parametrize("op", ["sum", "avg"])
@pytest.mark.parametrize("dtype", ["bfloat16", "fp8_e4m3", "fp8_e5m2"])
def test_ReduceScatterGinSdmaNarrowingDtypes(request, dtype, op):

at 4 MiB/rank, which is the mid-band 48-CTA probe the ladder already runs sum and avg at. That is 6 new launches, so GIN_PYTEST_RS_HW_CASES goes from 15 to 21 and the comment above it now itemises all three groups.

All six pass on hardware. On 8x MI350X gfx950 under ROCm 7.13, reduce_scatter_perf -b 4194304 -e 4194304 -R 2 -D 3 -A 1 -c 1:

bfloat16   sum   Out of bounds values : 0 OK
bfloat16   avg   Out of bounds values : 0 OK
fp8_e4m3   sum   Out of bounds values : 0 OK
fp8_e4m3   avg   Out of bounds values : 0 OK
fp8_e5m2   sum   Out of bounds values : 0 OK
fp8_e5m2   avg   Out of bounds values : 0 OK

That was worth running before adding them rather than after: fp8 through a pairwise fold that renarrows on every step is exactly the shape where a tolerance problem would have shown up, and it would have landed as a CI failure rather than a review comment.

I did not add a -d to the rccl-gin-rocshmem-sdma-rs-d3 entry. That one is the perf probe and adding dtypes there multiplies bench time without checking anything the pytest does not now cover, but say the word if you would rather have the coverage in both places.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

Low

These eleven lines reproduce gin_devtime::resolveLoopSkip (gin_sdma_devtime.h:34-46) statement for statement, including both skip clamps, and gin_sdma_devtime.h is already included at line 20. all_gather.cu:332 calls the helper. A later change to the shared tier ladder would silently miss ReduceScatter.

Suggested change

	}
	int loop = 0;
	int skip = 0;
	gin_devtime::resolveLoopSkip(totalBytesCta, loop, skip);

[dlamd1dai]

Agreed, and taken verbatim in 8efddab141. The eleven lines were a statement-for-statement copy of gin_devtime::resolveLoopSkip, including both skip clamps, from a header already included at line 20, and your point about the shared tier ladder drifting is the real risk — all_gather.cu:332 calls the helper, so ReduceScatter was the one caller that would have silently kept the old ladder.

int loop = 0;
int skip = 0;
gin_devtime::resolveLoopSkip(totalBytesCta, loop, skip);

The parameter is named perPeerBytes in the helper and the value passed is totalBytesCta; the thresholds compare against whatever is passed, so this is behaviour-identical, but the naming mismatch is worth being aware of if the helper ever grows a per-peer-specific clamp.

Exercised on hardware rather than just compiled, since this block only runs under --device_timing and would otherwise have shipped unverified:

#[rs-devtime] size      4194304 B  ctas 32  loop 10 skip 10  devtime      12.46 us  algbw   336.62 GB/s  busbw   294.54 GB/s

4 MiB total is under the 8 MiB mid-band threshold, so loop 10 skip 10 is the default arm resolving correctly.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

Low

These two lines restate the algBw/busBw formula that this PR's own gin_sdma_reducescatter::bandwidthGBps (gin_sdma_reducescatter_policy.h:184) computes and that ReduceScatterGetBw was rewritten to call at line 137, so the printed number is the one copy ReduceScatterPolicyBandwidth does not cover. all_gather.cu:362 calls its helper for exactly this augment line.

Suggested change

	double busBw = algBw * ((double)(nRanksGlobal - 1) / (double)nRanksGlobal);
	double algBw = 0.0, busBw = 0.0;
	gin_sdma_reducescatter::bandwidthGBps(count, (int)wordSize(type), sec, nRanksGlobal, &algBw, &busBw);

[dlamd1dai]

Agreed, and taken verbatim in 8efddab141.

double algBw = 0.0, busBw = 0.0;
gin_sdma_reducescatter::bandwidthGBps(count, (int)wordSize(type), sec, nRanksGlobal, &algBw, &busBw);

The framing in your note is what makes it worth doing: ReduceScatterGetBw was already rewritten to call the helper at line 137 and ReduceScatterPolicyBandwidth covers it, so this augment line was the single remaining copy and therefore the one place the formula could drift unnoticed. all_gather.cu:362 does the same thing for its own augment line, so this also matches the sibling.

totalBytes is still computed just above because the snprintf prints it, so that line is unchanged.

Verified on hardware through --device_timing, same run as the resolveLoopSkip thread:

#[rs-devtime] size      4194304 B  ctas 32  loop 10 skip 10  devtime      12.46 us  algbw   336.62 GB/s  busbw   294.54 GB/s

294.54 / 336.62 = 0.875 = 7/8, which is the (nranks-1)/nranks correction at 8 ranks, so the helper reproduces what the inline pair produced.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

Low

This cites parseAllGatherCtasEnv as precedent for leaving the host helpers unguarded, but gin_sdma_allgather_policy.h:158 opens #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) with that function inside it at :171, and all_gather.cu:148-155 works around the guard with its own device branch. The unguarded choice here may well be the better one; it is the precedent that is wrong.

Suggested change

	// gin_sdma_allgather_policy.h's parseAllGatherCtasEnv relies on.
	// arch. Plain `inline` is already host-only. Note gin_sdma_allgather_policy.h
	// does the opposite: it guards parseAllGatherCtasEnv, and all_gather.cu works
	// around that guard with a __HIP_DEVICE_COMPILE__ branch.

[dlamd1dai]

Good catch, and taken in 8efddab141. Tracing the citation through: gin_sdma_allgather_policy.h:158 does open #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) with parseAllGatherCtasEnv inside it at :171, and all_gather.cu:148-155 does carry a __HIP_DEVICE_COMPILE__ branch to work around that guard. So the comment's factual claim about the precedent was simply wrong.

The reasoning above it still stands on its own — HIP parses the whole translation unit in the device pass including bodies it will not codegen, so guarding these would break name lookup in ReduceScatterParseCtasEnv — it just should not have been attributed to a precedent that does the opposite. The comment now says so explicitly, which is more useful than silence: someone comparing the two headers will otherwise assume one of them is a mistake.

Thanks for checking the citation and not just the conclusion — a wrong attribution wrapped in correct reasoning is exactly the kind of thing that survives review, because it reads plausible.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

Low

bench_timeout is decided twice: line 215 binds it to GIN_PYTEST_TIMEOUT for every pytest kind, and this line re-binds it for ReduceScatter only. Broadcast is correct only because its outer_timeout at line 247 happens to equal line 215's value. A third pytest file that sets outer_timeout and misses this line gets apply_gin_pytest_inner_budget sizing against its own cap while timeout enforces Broadcast's, inverting the invariant lines 195-203 protect. Could this binding move to just after the esac?

[dlamd1dai]

Agreed, and moved as suggested in 8efddab141.

Your description of the failure mode is what makes this worth fixing now rather than later: the bug is latent, not active. Broadcast happens to be correct only because its outer_timeout equals the value line 215 binds, so today both files agree by coincidence. A third pytest file with its own outer_timeout that forgot this line would have apply_gin_pytest_inner_budget sizing its inner budget against its own cap while timeout enforced Broadcast's — which inverts exactly the invariant lines 195-203 exist to protect, and would surface as an orphaned mpirun session rather than a clean failure.

The binding now happens once, immediately after the esac, so it is structurally impossible for a file to set outer_timeout and miss it:

    esac
    # Bind once, after the esac, so every pytest file enforces the same cap it
    # sizes its inner budget against. Per-file re-binding inside the case would
    # silently pair one file's outer_timeout with another's `timeout`.
    bench_timeout="${outer_timeout}"

The per-file line in the ReduceScatter arm is gone. Behaviour for both existing files is unchanged, since each now gets its own outer_timeout where before Broadcast got the identical value by a different route.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

Low

Nothing in the PR sets NCCL_GIN_ANVIL_RS_CTAS, so the clamp composed here is only covered as arithmetic: gin_sdma_reducescatter_policy_test.cpp:166 pins the -V 8 plus RS_CTAS=64 regression, but swapping this call for a bare reduceScatterCtas(...) would reintroduce that out-of-bounds lsaBarrier index and every test would still pass. One pytest case carrying the pin through RCCL_TESTS_RS_XENV closes it, though it costs a launch and a GIN_PYTEST_RS_HW_CASES bump, so it may read better as a follow-up.

[dlamd1dai]

Agreed, and I am taking your suggestion that this reads better as a follow-up — so this one is not addressed in 8efddab141.

Your analysis is right and worth restating because the gap is subtle: gin_sdma_reducescatter_policy_test.cpp:166 pins -V 8 plus RS_CTAS=64 as arithmetic, but nothing in the PR ever sets NCCL_GIN_ANVIL_RS_CTAS against a real launch. Swapping the composed clamp for a bare reduceScatterCtas(...) would reintroduce the out-of-bounds lsaBarrier index and the whole suite would still pass, because the only thing testing the clamp is a unit test of the clamp.

I am deferring it rather than bundling it for two reasons. It costs a launch and another GIN_PYTEST_RS_HW_CASES bump, and this revision already raises that count from 15 to 21 for the narrowing-dtype coverage; I would rather land that increase, see the job's wall-clock, and then decide the budget for this one. Second, the pin belongs with the --device_timing smoke case you raised on the sibling thread — both are "wire an existing untested path through RCCL_TESTS_RS_XENV", and they share the plumbing.

I will file them together. If you would rather have this one in the current PR, say so and I will add it now; it is a small change, the hesitation is purely about CI budget.

[amd-argus]

Low

ReduceScatterDeviceTime and GinReduceScatterTimedKernel at line 369 are reachable only through --device_timing, which neither the pytest nor any gin-tests.json entry passes, so this whole block ships unexecuted. --devtime_check 1 already validates a timed kernel's output through the normal datacheck and returns an error (common.cu:1217-1228), and run-device-api-ci.sh runs an AllToAll devtime smoke on the same pattern. Worth one case as a follow-up?

[dlamd1dai]

Yes, worth a case, and I am taking it as the follow-up you propose rather than folding it in here — so this is not addressed in 8efddab141.

You are right that the block ships unexecuted. I should record that it became partly untrue by accident this round: I ran --device_timing 1 --devtime_check 1 by hand on 8x MI350X to validate the two shared-helper dedups raised on the other threads, since both resolveLoopSkip and bandwidthGBps live inside ReduceScatterDeviceTime:

#[rs-devtime] size      4194304 B  ctas 32  loop 10 skip 10  devtime      12.46 us  algbw   336.62 GB/s  busbw   294.54 GB/s
# Out of bounds values : 0 OK

So GinReduceScatterTimedKernel does run correctly and --devtime_check 1 does validate its output through the normal datacheck, exactly as you described from common.cu:1217-1228. But that was a manual run on my machine, which is not coverage — nothing in gin-tests.json or the pytest passes --device_timing, so the path is still one refactor away from silently breaking, and the fact that I had to reach for it by hand rather than point at an existing case rather makes your argument.

run-device-api-ci.sh running an AllToAll devtime smoke on the same pattern is a good template to copy. I will file this together with the NCCL_GIN_ANVIL_RS_CTAS pin from the sibling thread, since both are the same shape — push an existing untested path through RCCL_TESTS_RS_XENV and account for the extra launch in GIN_PYTEST_RS_HW_CASES.

[amd-argus]

Low

RS_NP here, and RS_TIMEOUT_S and RS_CONN_RETRIES at lines 82-83, parse env vars with a bare int() at module scope, so a value like 900s raises at import and pytest reports a collection error for the whole file instead of skipping. test_Broadcast.py:158 already has _env_int(name, default) for exactly this, used for the BCAST vars at lines 169, 177 and 178. Could these three use it?

[dlamd1dai]

Agreed, and fixed in 8efddab141 using the existing helper as you suggested.

The failure mode is worse than a bad value being ignored: a bare int() at module scope raises during import, so pytest reports a collection error for the entire file rather than skipping or falling back. A value like 900s — plausible given GIN_PYTEST_RS_TIMEOUT is itself written with the s suffix — would take out every ReduceScatter test rather than one.

test_Broadcast.py:158 already had exactly the right helper, so I copied it rather than inventing a variant, and all three vars now use it:

RS_NP = _env_int("RCCL_TESTS_RS_NP", 0)
RS_TIMEOUT_S = _env_int("RCCL_TESTS_RS_TIMEOUT_S", 900)
RS_CONN_RETRIES = _env_int("RCCL_TESTS_RS_CONN_RETRIES", 5)

RS_NP keeps its existing <= 0 fallback chain (ROCR_VISIBLE_DEVICES, then HIP_VISIBLE_DEVICES, then rocminfo) underneath, so a bad value now lands on device autodetection instead of aborting collection. The int() inside that rocminfo fallback is left as-is, since it parses command output rather than an env var and is already inside a try/except.

Longer term the duplication between the two files is worth pulling into a shared conftest helper, but that is a wider change than this PR and I would rather not drag test_Broadcast.py into it here.

[amd-argus]

Resolved by 8efddab141.

[amd-argus]

ARGUS

This PR changes 1,750 lines across 14 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 relaxed lsaBar.sync at reduce_scatter.cu:244, still open from the previous round.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

[amd-argus]

Medium

gin_sdma_reduce::combine, preOp and postOp have no caller outside reduce_scatter.cu, so the bit-match-with-verifiable.cu contract stated at the top of this header is checked only by the -D 3 datacheck on an 8-GPU GIN node, which run-gin-ci.sh runs non-gating. Reverting the half/bf16 narrow-on-every-step to a float accumulator would leave every gtest green. Could a kernel in gin_sdma_reducescatter_gpu_test.cpp fold these ops over fixed inputs per (op, T) against expected values written out in the test?

[dlamd1dai]

Good catch — fixed in f360676. Of the three findings in this round, this is the one with teeth.

ReduceOpsNarrowEveryStep in gin_sdma_reducescatter_gpu_test.cpp folds {B, 1, 1, 1, 1, 1, 1, 1} over eight ranks for float, int32, half, bfloat16, fp8_e4m3 and fp8_e5m2 under sum/prod/max/min/avg, against expected values written out in the test as you suggested.

The choice of inputs is what makes it non-vacuous. B is picked so that one ulp of B exceeds 2, which puts B+1 strictly inside the ulp — never on a tie, so the result does not depend on ties-to-even — and therefore pins the step-wise sum at B. A float accumulator would instead reach B+7 and narrow once, landing an ulp or two higher. float and int32 are the controls: neither narrows, so they see the full B+7 and stay green under either implementation.

T	B	step-wise sum / avg	float accumulator
half	4096	4096 / 512	4104 / 513
bfloat16	512	512 / 64	520 / 65
fp8_e4m3	32	32 / 4	40 / 5
fp8_e5m2	32	32 / 4	40 / 5

I computed the expected column independently before building, rather than reading it back off a run, so the test is asserting a derived contract and not just pinning current behaviour. On gfx950 all five ReduceScatterGpu cases pass. As a negative control I rebuilt the same test with the fold accumulating in float and narrowing once: it fails exactly those eight low-precision sum and avg cases at exactly the predicted values, and float/int32 stay green.

The test also pins ncclSum=0 .. ncclAvg=4, since combine switches on the raw int while SPECIALIZE_REDUCE_KERNEL hands it an ncclRedOp_t. fp8 prod is left out because that macro never dispatches it there.

Two things worth flagging on the plumbing. The header asks to be included after common.h, but common.h is rccl-tests' MPI harness and this target has no business linking it, so the test reproduces only what the header actually reads: rccl.h (for NCCL_VERSION_CODE, ncclRedOp_t and __half, which is why the target gains ${PROJECT_BINARY_DIR}/include, mirroring RCCL_COMMON_INCLUDE_DIRS at test/CMakeLists.txt:163), the HAVE_BF16/HAVE_FP8 guards, rccl_float8.h, and ENABLE_DEVICE_API — opened locally, since this is pure arithmetic and the coverage should not blink out depending on how librccl was configured.

And rccl_float8.h typedefs different underlying types per arch: fnuz on gfx942, IEEE on gfx950, and a software struct on everything else. 32, 4 and 1 are exact in all three and the mantissa widths and RNE behaviour are identical, so the expectations hold across them. I compile-checked gfx90a, gfx942 and gfx1100 to cover the fallback and fnuz paths; only gfx950 was available to run.

[amd-argus]

Resolved by f3606768a3.

[amd-argus]

Low

sliceBaseCount, sliceBytes, reduceScatterSliceOffset and pickSdmaThreshold are character-for-character copies of chunkBaseCount, chunkBytes, allGatherRecvSliceOffset and pickSdmaThreshold in gin_sdma_allgather_policy.h, and gin_sdma_broadcast_policy.h is a third copy of the same family. Divergence between copies is silent: the ReduceScatter unit test locks one behaviour while AllGather drifts, and nothing goes red. Consolidating means editing AllGather, so this reads as a follow-up rather than something to fix in this PR.

[dlamd1dai]

Agreed, and your framing is right: this belongs in a follow-up rather than here.

The duplication is real: sliceBaseCount, sliceBytes, the slice-offset helper and pickSdmaThreshold are the AllGather versions with the names changed, and Broadcast makes a third copy. But consolidating means editing gin_sdma_allgather_policy.h and gin_sdma_broadcast_policy.h, which drags all_gather.cu and broadcast.cu into a PR whose hardware coverage is entirely ReduceScatter — I would be changing the two collectives I am least able to regression-test in this branch.

I will file it together with the other follow-ups from this round (the NCCL_GIN_ANVIL_RS_CTAS pytest pin and the --device_timing smoke case). The shape I have in mind is a single gin_sdma_slice_policy.h holding the count/bytes/offset arithmetic and the threshold pick, with the three per-collective headers reduced to their CTA ladders, which are the part that genuinely differs.

[amd-argus]

Low

test_rs_pytest_kind_preserves_args can never fail on its own: it differs from test_pytest_kind_preserves_args only in the name/bin strings and a one- instead of two-entry env list, and parse_gin_config.py branches on neither bin nor kind (parse_config at parse_gin_config.py:49 copies both through, format_rows at :78 joins them positionally). Could the existing case be parametrized over both filenames instead?

[dlamd1dai]

Good catch — fixed in f360676.

Tracing it through: parse_config copies bin and kind straight through, and format_rows joins the fields positionally. Neither branches on either value, so there is no code path test_rs_pytest_kind_preserves_args could reach that test_pytest_kind_preserves_args did not — it could only ever have failed as a duplicate.

Parametrized as you suggested. Since the file is unittest rather than pytest, the two filenames now run through a subTest loop in the single case. I kept the differing env list lengths (two entries and one) so the -x-prefixed join is still pinned at both arities, and added an assertion on parts[1] so the name field is covered too.

[amd-argus]

Resolved by f3606768a3.

[amd-argus]

ARGUS

This PR changes 1,984 lines across 14 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: Looks approval ready. One open item if you want it here: the avg skip added at reduce_scatter.cu:572 has no test that can fail on it, since all 21 pytest launches assert exit status only.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

[amd-argus]

Medium

The deviceImpl == 0 && qualifier added here is not covered by anything that can fail on it. Dropping it makes the four avg launches (float, bfloat16, fp8_e4m3, fp8_e5m2) continue past their only op x type pair, return testSuccess and exit 0, and the pytest asserts exit status alone at test_ReduceScatterGinSdma.py:201. Could the skip matrix become a pure predicate in gin_sdma_reducescatter_policy.h, pinned by the policy UT that already runs as rccl-gin-rs-policy-ut?

[dlamd1dai]

Superseded by #12747 and closing in favour of it. Same commits, same content; only the head branch moved.

Why. This PR was opened from a fork, and the required RCCL PR gate check can never pass that way. Fork pull requests get no secrets, so the RCCL coco PR workflow skips its Setup job, the coco smoke never runs, and the gate fails closed within seconds. Clearing it here would need a maintainer to dispatch the workflow by hand and then override the check, on every push. #12747 is branched in ROCm/rocm-systems, where Setup succeeds and coco dispatches to both the Ruby and OCI clusters on its own.

Please keep reading this thread. 14 of the 18 review threads here are still unresolved, and inline conversations do not carry across to a new pull request. Argus will re-review #12747 and should re-surface anything still unfixed, but the replies and reasoning written here only exist on this PR, so check them before merging the replacement.

No further work is needed on this branch. Closing rather than leaving it open so it does not keep consuming CI or show up as a duplicate for reviewers; it stays readable and can be reopened.