test(rccl): extend rccl_wrap.cc host coverage to 24 more functions (52.6% lines) by aduforest · Pull Request #10909 · ROCm/rocm-systems

aduforest · 2026-08-28T19:15:04Z

JIRA ID : AICOMRCCL-2195

What

Extends host-only microtest coverage for src/rccl_wrap.cc, stacked on
#10841. Test-only — src/rccl_wrap.cc is untouched.

Whole-file coverage: 16.43% → 52.60% lines, 40.48% → 83.33% functions,
8.53% → 34.06% branches.

24 more functions covered, in six natural clusters:

Cluster                              Functions
Trivial leaf                         symkHostRedOpToDev
Env/override                         rcclUpdateCollectiveProtocol, rcclUpdateThreadThreshold,
                                      rcclOverrideProtocol, rcclOverrideAlgorithm, plus the
                                      remaining getenv-driven paths of rcclSetPxn /
                                      rcclSetP2pNetChunkSize (cached fast-path covered in #10841)
Direct/hierarchical eligibility      rcclUseHierarchicalAllGather, rcclUseAllGatherDirect,
                                      rcclUseReduceScatterDirect, rcclUseHierarchicalReduceScatter
Tuning knobs                         rcclOptThreadBlockSize, commSetUnrollFactor,
                                      rcclCommSetP2pShiftSize, rcclOverrideChannels
CE / DDA gates                       rcclUseCeAllReduce, rcclDdaEnabled
Symmetric-kernel low-tier arms       rcclSymkQuery, rcclSymKGetInfo, getFirmwareVersion

Not covered, deferred to a future PR: rcclSetPipelining; the 10-function
WarpSpeed group (not even compiled into this binary — ENABLE_WARP_SPEED is
off); and the High-tier group that shares the DDA/CE/symmetric-kernel
abort-floor surface — rcclSelectAllReduce/AllGather/ReduceScatter,
rcclHierarchicalAlgoInfo, rcclGetAlgoInfo, rcclGetCollImplInfo, and the
deep (post-guard) path of rcclSymkQuery/rcclSymKGetInfo. All of these
need several abort() floors in wrap_stubs.cc upgraded to controllable
hooks first — a larger, separate piece of work, split out per the same
precedent that separated this PR from #10841.

How

Five small seams were added first (fakes/wrap_stubs.cc/.h), all
following patterns already established elsewhere in this test suite
(init-test.cc/fakes/nccl_fakes.cc):

RCCL_PARAM redirector upgraded from a fixed compile-time default to a
g_loadParam std::function hook — same mechanism and "RCCL_" + env
convention as init-test.cc's redirect — so a test can flip one param's
value between cases.
ncclGetEnv and rccl_wrap.cc's several bare getenv() call sites (a
macro can't intercept those) now route through a shared settable map,
following fakes/init_fakes.cc's micro_getenv/SetMicroEnv/
link-level-getenv-override pattern, self-contained in this file.
ncclGroupDepth, rcclUseAinic(), ncclPxnDisable() added as simple
stubs matching their real common-case defaults.
amd_smi_getFirmwareVersion upgraded from an abort floor to a settable
hook (same shape as fakes/nccl_fakes.cc's), for getFirmwareVersion.

Each seam was verified to preserve prior behavior exactly (all existing
tests still passed immediately after adding it, before any new test used
it), and, where applicable, that breaking it produces a real, observable
failure.

Several functions in rccl_wrap.cc have real defaults that already favor
the "enabled" path (e.g. rcclParamHierarchicalAllGather, rcclDdaEnable,
rcclParamChannelTuningEnable), so their main branches were reachable
without the seam — it was only needed to reach the disabled arm. Others
(rcclParamHierarchicalReduceScatter, rcclUseCeAllReduce's
rcclParamCeAllReduce) default to fully blocking the function, so every
branch needed the seam.

Seams (process-isolation note)

rcclSymkQuery's four early-return guards are tested via
RUN_ISOLATED_TEST, not plain TEST(): every guard is one mutation away
from falling through into the ncclSymkInitOnce()/etc. abort floor, which
would crash the whole binary in an in-process test instead of failing one
case. Confirmed this is a real risk, not theoretical, while mutation-testing
the null-comm guard — a deliberately-broken version segfaults cleanly
inside its own isolated process, caught by 2 of 4 tests, with no impact on
the rest of the run.

Mutation

Every new function had at least one mutation applied directly against
src/rccl_wrap.cc, rebuilt, confirmed the right test failed, then reverted.
git diff --stat -- src/ confirmed empty after every revert.

Two genuine equivalent mutants found and documented rather than forced:

Residual	Why it is accepted
`rcclOverrideProtocol`'s `protoVal > NCCL_PROTO_UNDEF` guard mutated to `>=`	`protoVal` only ever reaches this line as a successfully-parsed index (always > -1) or after an early return on parse failure — it can never actually equal `NCCL_PROTO_UNDEF` here, so no input distinguishes `>` from `>=`.
`rcclSymKGetInfo`'s null-arg check with one arm removed (e.g. `algo == nullptr`)	`rcclGetCollImplInfo` (what this function falls through to whenever `rcclSymkQuery` returns false) performs the identical `algo`/`protocol`/`maxChannels` null check redundantly, so no input can observe a difference.

Three boundary mutations initially survived because the existing test value
wasn't at the exact boundary (rcclOptThreadBlockSize's max-threads clamp,
rcclCommSetP2pShiftSize's bit-reversal gate, rcclOverrideChannels' range
match); each was fixed by moving the test value to the boundary itself. Two
threshold-swap mutations (rcclSetPxn/rcclSetP2pNetChunkSize's
arch-threshold selection) similarly survived because the existing boundary
values happened to classify the same way under the swapped thresholds; fixed
by adding a value strictly between the two archs' real thresholds.

Test plan

Binary	Result
`rccl-UnitTestsMicroWrap`	136/136
Default order	pass
`--gtest_shuffle`	pass
`--gtest_shuffle --gtest_repeat=5`	pass
`--gtest_shuffle --gtest_repeat=10 --gtest_random_seed=42`	pass
Coverage (`llvm-cov report`)	52.60% lines / 83.33% functions / 34.06% branches on `rccl_wrap.cc`
Fresh clean rebuild (build dir wiped)	pass, no new warnings

No CMake change beyond a documentation-accuracy fix to an existing comment.
Built host-only against the hipified snapshot — no GPU, no librccl.so, no
HIP runtime.

therock-pr-bot · 2026-08-28T19:32:41Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/rccl/test/host/fakes/wrap_stubs.cc`, `projects/rccl/test/host/fakes/wrap_stubs.h`, `projects/rccl/test/host/wrap-test.cc`; no test file found
🚫 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?

Reviewer feedback: don't add a new binary unless necessary. The original split only ruled out the init binary (real duplicate-symbol conflict on commSetUnrollFactor/rcclCommSetP2pShiftSize); combining with the p2p binary was never actually checked and turns out to work. Merged wrap-test.cc/fakes/wrap_stubs.cc into rccl-UnitTestsMicro, deduped two symbol collisions found along the way (shared logging globals, and a stubbed IsArchMatch vs. the real one), and deleted the now-unused rccl-UnitTestsMicroWrap binary and its categories yaml. 95/95 tests pass (default order, shuffle, shuffled x10). Coverage on rccl_wrap.cc unchanged. Test-only -- src/rccl_wrap.cc untouched.

… commit One comment claimed wrap-test.cc already raises ncclDebugLevel/Mask via ScopedDebugLogging before capturing INFO logs -- not true yet on this branch (no INFO-gated test exists here until a later batch); reworded to say why the differing mask default isn't observable today instead. Also fixed comment-column alignment in CMakeLists.txt's two source lists. No functional change.

…ing link stubs Prerequisite infrastructure for the next coverage batch (no new tests yet; tracked separately). Adds five small, precedented seams so wrap-test.cc can reach functions that need something other than a plain comm/topology setup: - RCCL_PARAM redirector upgraded from a fixed deftVal to a g_loadParam std::function hook, same mechanism and "RCCL_" + env convention as init-test.cc's redirect, replicated here rather than linking fakes/nccl_fakes.cc (this file's fakes stay self-contained). - ncclGetEnv and rccl_wrap.cc's several bare getenv() call sites (which a macro can't intercept) now route through a shared settable map, following fakes/init_fakes.cc's micro_getenv/SetMicroEnv/link-level-getenv-override pattern (new fakes/wrap_stubs.h exposes the test-control API). - ncclGroupDepth (was undefined -- link only worked because --gc-sections stripped it from code paths not yet reached), rcclUseAinic(), and ncclPxnDisable() added as simple stubs matching their real common-case defaults. Verified each seam preserves prior behavior exactly (49/49 existing tests still pass) and, where applicable, that breaking it produces a real, observable failure (temporarily flipped each new stub/hook and confirmed a test would have caught a regression, then reverted). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

…uster Adds 27 tests covering 8 more rccl_wrap.cc functions: - symkHostRedOpToDev: pure switch, no setup needed. - rcclUpdateCollectiveProtocol: user-override gate, one arch/size LL threshold arm, gfx120x delegation + NCCL_P2P_DISABLE override (exercises the new ncclGetEnv seam), and the multi-node minMaxLLRange-driven arm including its warn-once undefined-tuning fallback. The gfx950/gfx942 ReduceScatter arms (same shape, different constants) and the ENABLE_LL128 nested arm (off in this build) are out of scope, documented at the test site. - rcclUpdateThreadThreshold: tuned-value and user-override arms. - rcclOverrideProtocol / rcclOverrideAlgorithm: unset-passthrough, successful-override, ignored-combo, and unmatched-string arms. rcclOverrideAlgorithm is structurally identical, so only its unset-passthrough and successful-override arms are re-verified. - rcclSetPxn / rcclSetP2pNetChunkSize: closes the previously-documented getenv-driven-path gap from an earlier batch (the cached fast-path was already covered). rcclSetP2pNetChunkSize's final unsupported-arch WARN arm is dead code (the guard above it already returns early whenever neither arch matches) -- documented, not contrived. Mutation-tested every function above (one applied at a time against src/rccl_wrap.cc, rebuilt, confirmed the right test failed, reverted). Two mutations initially survived and were fixed by strengthening the test rather than accepting a false pass: - rcclSetPxn/rcclSetP2pNetChunkSize's arch-threshold swaps survived because the existing boundary values (64, 31, 16) happened to classify the same way under both the real and swapped thresholds; added nRanks values strictly between the two archs' real thresholds to actually distinguish them. One mutation (rcclOverrideProtocol's `protoVal > NCCL_PROTO_UNDEF` guard, mutated to `>=`) is a genuine equivalent mutant: protoVal can never equal NCCL_PROTO_UNDEF at that point in any reachable execution, so no input could ever distinguish the two. Documented at the test site rather than forced. 74/74 tests pass. Test-only; src/rccl_wrap.cc untouched. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Adds 19 tests covering 4 more rccl_wrap.cc functions: - rcclUseHierarchicalAllGather: nNodes gate, initialization gate, threshold boundary (both arms), and the param-disabled arm (via the g_loadParam seam). Its real default (1, "enabled") makes the interesting branches reachable without the seam, unlike its ReduceScatter sibling below. - rcclUseAllGatherDirect: param-disabled gate, auto-threshold selection on gfx950, the CTA-policy-ZERO single-node disable, and the rank-multiple- of-8 gate. - rcclUseReduceScatterDirect: param-disabled gate, non-gfx950 gate, and the per-node-count threshold table (2/4/8/16 nodes). The PXN-disabled early return is out of scope: ncclPxnDisable is a plain stub always returning 0, not yet a controllable seam -- deferred, documented at the test site. - rcclUseHierarchicalReduceScatter: the always-false default (its real param default is 0, unlike AllGather's 1) plus the enabled/in-threshold arm via the seam. Mutation-tested every function (one applied at a time against src/rccl_wrap.cc, rebuilt, confirmed the right test failed, reverted). Found and fixed three real gaps this way rather than accepting false passes: - rcclUseHierarchicalAllGather's threshold boundary needed an exact-match test value; the existing one (1024 bytes against a 32MiB threshold) was nowhere near the boundary and couldn't distinguish <= from <. - rcclUseReduceScatterDirect's 4-node branch had no test at all -- added one, which then caught a real threshold-constant mutation. - rcclUseHierarchicalReduceScatter's nNodes>=8 gate needed an 8-node case specifically: at 8 nodes its own temp-buffer threshold is already nonzero (64MiB), which is what actually distinguishes the real ">= 8" gate from a mutated ">= 16" one -- the existing 16-node test couldn't. 93/93 tests pass. Test-only; src/rccl_wrap.cc untouched. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Adds 19 tests covering 4 more rccl_wrap.cc functions: - rcclOptThreadBlockSize: same isolated-static-cache shape as the already- covered rcclSetDefaultBuffSizes. Covers the user-override path (used directly, warp-aligned; rounded up when not aligned; clamped to max), the TREE-algorithm default, the single-node default, and the ReduceScatter small-count optimization. - commSetUnrollFactor: valid/invalid user-override, and four of the arch- based default branches (gfx950 single- and multi-node, gfx908, an unlisted arch). The "not built for this arch" fallback arms (both the override path's and the default path's) are unreachable: wrap_stubs.cc's ncclDevFuncUnrollGenerated is a hardcoded all-true array, not yet a settable hook -- deferred, documented at the test site. - rcclCommSetP2pShiftSize: both branches of the bit-reversal-vs-exact-value guard. Not cached (the RCCL_PARAM redirector doesn't cache), so no isolation needed. - rcclOverrideChannels: both early-return gates (too few nodes, single GPU per node), the param-disabled gate, and both outcomes of the per-entry threshold match (channel count applied when within CTA bounds, left untouched when conflicting). Mutation-tested every function (one applied at a time against src/rccl_wrap.cc, rebuilt, confirmed the right test failed, reverted). Found and fixed four real gaps this way: - Three boundary mutations (rcclOptThreadBlockSize's max-threads clamp, rcclCommSetP2pShiftSize's bit-reversal gate, rcclOverrideChannels' range match) initially survived because the existing test values weren't at the exact boundary; adjusted each to the boundary value itself. One of these (the clamp guard) turned out to be unobservable in nThreads alone right at the boundary -- clamping a value to itself doesn't change the visible result -- so a different, cleanly-observable mutation (the warp-alignment rounding formula) was used instead. - rcclOverrideChannels' matched-threshold test never set comm->nChannels, so maxNChannels silently clamped to 0 and the override branch was never actually reached -- caught by the test itself failing before any mutation was even applied, fixed by setting a realistic channel count. 112/112 tests pass. Test-only; src/rccl_wrap.cc untouched. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Adds 17 tests covering 2 more rccl_wrap.cc functions: - rcclUseCeAllReduce: rcclParamCeAllReduce's real default (0) fully blocked this function before the g_loadParam seam existed, so every branch needed the seam. Covers the disabled-by-default warn-once gate, the bias-buffer rejection, symmetric-support/single-node/count-divides- evenly guards, the message-size limit, the CTA-policy-ZERO requirement (and force bypassing it), unsupported ops, Float8 rejection, and the fully-valid success path. One test needed RcclUnitTesting:: ScopedDebugLogging: "CE AllReduce not enabled" is an INFO log (gated on ncclDebugLevel, which defaults to suppressed), unlike this file's other WARN-based messages which log unconditionally -- caught by the test itself failing before any mutation was applied. - rcclDdaEnabled: rcclParamDdaEnable's real default (1) and ncclGroupDepth's new stub (0) both favor the "enabled" path, so most branches are reachable without the seam; only needed for the explicitly-disabled case. Covers the gfx1250/gfx942/gfx950 arch branches (including the gfx942-vs-gfx950 default-threshold fallback), the gfx942 rank floor, and the unsupported-arch fallthrough. Mutation-tested every function (one applied at a time against src/rccl_wrap.cc, rebuilt, confirmed the right test failed, reverted). 129/129 tests pass. Test-only; src/rccl_wrap.cc untouched. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Adds a new settable hook and 9 tests covering 3 more rccl_wrap.cc functions: - rcclSymkQuery (static): only its four early-return guards (null comm, no symmetric support, unsupported collective, unmapped reduction op). Everything past ncclSymkInitOnce() is still an abort-floor stub, deferred to a future batch that upgrades the whole DDA/CE/symmetric-kernel abort-floor surface to controllable hooks. These tests run isolated, not in-process: a mutation that loosens any of the four guards would let execution fall through into that abort floor, which would crash the whole binary in an in-process test instead of just failing one case -- confirmed this is a real risk, not theoretical, while mutation-testing the null-comm guard (a deliberately-broken version segfaults cleanly inside its own isolated process; caught by 2 of 4 tests, no impact on the rest of the run). - rcclSymKGetInfo: only its null-arg-check arm; the success path calls rcclSymkQuery (fine) but the fall-through path calls rcclGetCollImplInfo, itself High-tier (same deferred surface). Mutation-testing found this guard is a genuine equivalent mutant: rcclGetCollImplInfo performs the identical algo/protocol/maxChannels null check redundantly on the fall-through path, so no input can observe a difference -- documented at the test site rather than forced. - getFirmwareVersion: needed amd_smi_getFirmwareVersion upgraded from an abort floor to a settable std::function hook (same shape as fakes/nccl_fakes.cc's), exposed via fakes/wrap_stubs.h. Covers both the success (reports the canned firmware version) and failure (-1) paths. Mutation-tested every function (one applied at a time against src/rccl_wrap.cc, rebuilt, confirmed the right test failed or documented why not, reverted). 136/136 tests pass. Test-only; src/rccl_wrap.cc untouched. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Two comments described the pre-batch state rather than the code as it now stands, found during a final review pass: - wrap_stubs.cc's file header still said RCCL_PARAM invocations return their compile-time default directly with nothing per-test-controllable; this batch's first commit replaced that with the g_loadParam redirector. - wrap-test.cc's file header and CMakeLists.txt's matching comment still said "17 low-dependency helpers", with rcclOverrideChannels named as out of scope -- both now covered. Reworded to describe what's NOT yet covered instead of enumerating what is, so this doesn't go stale again the next time a function gets added. Comment-only; no test or behavior change. 136/136 still passing. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Found during a follow-up audit for the same category of concern already raised on this file: wrap_stubs.cc's `thread_local int ncclGroupDepth = 0;` copies the real thread_local variable's initializer (group.cc:31), but unlike NCCL_NUM_ALGORITHMS etc. there's no importable constant to static_assert against -- it's a plain variable definition. Confirmed rcclUseAinic()/ncclPxnDisable() (the other two seams added earlier in this batch) do NOT have this issue: both real implementations are genuinely runtime/hardware-state-dependent, with no fixed default value being mirrored. Added GroupDepth_InitializerMatchesProductionSource, the same run-time tripwire as ParamDefaults_MatchProductionSource: reads the real group.cc source (new GROUP_CC_PATH CMake define, same pattern as CONNECT_CC_PATH/ ENQUEUE_CC_PATH) and confirms the exact initializer text is still there. Verified it actually fires by temporarily breaking the search string, confirming the test failed with the right message, then reverting. 137/137 tests pass. Test-only; src/rccl_wrap.cc untouched. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Surfaced only after rebasing onto the binary-merge commit: wrap-test.cc's own g_loadParam and fakes/nccl_fakes.cc's (p2p's, already part of this binary) are both non-static globals with the same name -- a duplicate symbol at link time. nccl_fakes.h already exports g_loadParam for exactly this kind of reuse (its own doc comment says as much); wrap-test.cc just hadn't needed it before since it used to build as its own separate binary. Dropped the local declaration and included fakes/nccl_fakes.h instead -- same instance, wrap-test.cc's own RCCL_PARAM macro still prepends "RCCL_" at its call site same as before, p2p's NCCL_PARAM macro still doesn't. 183/183 tests pass (default order, shuffled with several seeds, shuffled x10/x15 with repeats) -- no cross-suite leakage from sharing the hook. Coverage on rccl_wrap.cc unchanged: 52.60%/83.33%/34.06%. Test-only.

…linked in Rebasing onto current develop pulled in another team's group-test.cc, which compiles the real src/group.cc directly into this same binary -- and group.cc is exactly where the real ncclGroupDepth lives. That made wrap_stubs.cc's own hand-copied `thread_local int ncclGroupDepth = 0;` a genuine duplicate symbol at link time (confirmed: ld.lld rejected it pointing at both group.cc:32 and wrap_stubs.cc:175). Deleted the stub and, since there's no longer a copy to drift, the GroupDepth_InitializerMatchesProductionSource test that guarded it -- that test's only job was catching wrap_stubs.cc's copy going stale, and there's nothing left to go stale. Cleaned up the now-duplicate GROUP_CC_PATH definitions and stale comments left over from group-test.cc and this branch independently adding the same CMake plumbing. 186/186 tests pass (default order, shuffled with multiple seeds, shuffled x10/x15 with repeats). Coverage on rccl_wrap.cc unchanged: 52.60%/83.33%/34.06%. Test-only.

github-actions · 2026-08-29T02:13:58Z

RCCL Perf-Regression Gate: ⚠️ NO VERDICT (not measured)

This run did not produce a usable answer. It is neither a PASS nor a regression — treat the perf gate as not run for this change.

Why:

this report predates the trustworthiness flag, or a group could not be scored

Mode: detect (reference vs candidate)
Thresholds: small 17.2% · mid 12.3% · large 12.9%
Keys compared: 0 · Confirmed regressions: 0 · Inconclusive: 0
Provenance: n/a

Per-collective breakdown

group	keys	regressions	inconclusive
all_gather_perf-d=bfloat16-default	0	0	0
all_gather_perf-d=float-default	0	0	0
all_reduce_perf-d=bfloat16-default	0	0	0
all_reduce_perf-d=float-default	0	0	0
broadcast_perf-d=bfloat16-default	0	0	0
broadcast_perf-d=float-default	0	0	0
reduce_scatter_perf-d=bfloat16-default	0	0	0
reduce_scatter_perf-d=float-default	0	0	0

amd-argus

Automated Review Guard for Upfront Scrutiny

Findings

6 finding(s) are posted inline, on the lines they refer to.

projects/rccl/test/host/wrap-test.cc (Low) Branch coverage of 34.06% is just under the repo's own documented Basic acceptance bar of 35%
projects/rccl/test/host/wrap-test.cc (Low) Follow-up, untouched production code: size_t threshold compared against -1 at rccl_wrap.cc:1338 makes the clamp dead

Bottom line: No blocker. The one theme worth acting on before merge: four negative tests and five threshold arms still pass with the line they are named for deleted, and fixing those one-line gaps also clears the repo's own 35% branch bar.

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

amd-argus · 2026-08-30T06:54:18Z

+      []() {
+        SetMicroEnvAbsent("RCCL_DIRECT_ALLGATHER_THRESHOLD");
+        ncclComm* comm = MakeCommWithArch("gfx950");
+        comm->nNodes = 1;


UseAllGatherDirect_CtaPolicyZeroDisablesOnSingleNode still passes with the CTAPolicy guard at rccl_wrap.cc:739-743 deleted. MakeCommWithArch leaves nRanks = 1 (wrap-test.cc:91), so rankMultiple = 1 % 8 = 1 and the final !rankMultiple conjunct at rccl_wrap.cc:763 returns false regardless. Same shape at lines 1245, 1310 and 1321. Could we set comm->nRanks = 8 here, and comm->nNodes = 8 in the two ReduceScatter cases, the way the sibling test at line 1259 already does?

amd-argus · 2026-08-30T06:54:18Z

+        ncclComm* comm = MakeCommWithArch("gfx942");
+        comm->nNodes = 2; // avoid the nNodes==1 arm so the algorithm arm is what sets nThreads
+        ncclTaskColl info{};
+        info.algorithm = NCCL_ALGO_TREE;


OptThreadBlockSize_TreeAlgorithmUsesMaxThreads passes with the arm it names deleted. ncclTaskColl info{} leaves info.protocol at 0, which is NCCL_PROTO_LL, so the standalone if at rccl_wrap.cc:1632 is overwritten by the LL arm at rccl_wrap.cc:1638:

rccl_wrap.cc:1632 Tree arm nThreads = maxNthreads[SIMPLE] = RCCL_DEFAULT_MAX_NTHREADS = 256 rccl_wrap.cc:1638 LL arm nThreads = maxNthreads[LL] = RCCL_LL_MAX_NTHREADS = 256 (rccl_common.h:50-53: both constants are 256)

Adding info.protocol = NCCL_PROTO_SIMPLE; makes deleting rccl_wrap.cc:1632 leave nThreads at -1. OptThreadBlockSize_SingleNodeUsesHalfThreads at line 1517 needs the same pin.

amd-argus · 2026-08-30T06:54:18Z

+      "Wrap_OptThreadBlockSize_UserOverrideUsedDirectlyWhenAligned",
+      []() {
+        g_loadParam = [](const char* env, int64_t deft) {
+          return std::strcmp(env, "RCCL_THREADS_PER_BLOCK") == 0 ? int64_t(192) : deft; // 192 = 3*64, aligned


Five threshold arms are never entered by any new test, so each one survives deletion. The min-threads bump at rccl_wrap.cc:1626-1629 is only ever reached with nThreads already at 192, so its body never runs; the gfx942 1 << 17 arm at rccl_wrap.cc:1404 is untested because the only gfx942 case uses nRanks = 64; the 2-node upper bound at rccl_wrap.cc:1350 is never probed; the nNodes == 16 disjunct at rccl_wrap.cc:1353 is unreachable behind the hard-limit return at rccl_wrap.cc:1348; and the second and third env probes at rccl_wrap.cc:340 and rccl_wrap.cc:343 are never the one that fires. The other four anchor at lines 1133, 1334, 1382 and 903.

amd-argus · 2026-08-30T06:54:18Z

+// ===========================================================================
+
+TEST(WrapMicrotest, SetUnrollFactor_ValidUserOverrideSucceeds) {
+  g_loadParam = [](const char* env, int64_t deft) {


g_loadParam is installed by bare assignment here and restored by hand at line 1559, so any path that leaves the body early, an ASSERT_* added later or a throw out of CaptureLog, leaks the hook into every subsequent test in the binary. test/host/ScopedHook.h exists for exactly this and p2p-test.cc:1644 already uses it on this same global. Could we use it at all seven sites, lines 1553, 1564, 1610, 1623 and 1902 for g_loadParam plus 2046 and 2055 for g_amdSmiGetFirmwareVersion? That also removes the hand-written copies of the stub defaults at lines 2051 and 2057.

amd-argus · 2026-08-30T06:54:18Z

+
+// Returns a pointer into the map's own std::string: re-scripting a name
+// invalidates a pointer a caller may still hold.
+const char* micro_getenv(const char* name);


micro_getenv is declared here but has no caller anywhere outside wrap_stubs.cc, where ncclGetEnv at wrap_stubs.cc:165 is its only user and sits in the same TU. SetMicroEnv, SetMicroEnvAbsent, ClearMicroEnv and g_amdSmiGetFirmwareVersion all have real wrap-test.cc callers. Could it move into the anonymous namespace rather than being part of this header's API?

amd-argus · 2026-08-30T06:54:18Z

+  comm->config.maxCTAs = 64;
+  // bytesPerRank = divUp(nBytes, nRanks) = divUp(16384, 8) = 2048, exactly at
+  // maxByteThreshold -- distinguishes the range check's <= from a plain <.
+  comm->minMaxChannelThresholds[RCCL_AR_TUNABLE][0][0] = 1;    // minByteThreshold (nonzero: 0 collides with CHAN_THRESHOLDS_UNDEFINED)


Suggested change

comm->minMaxChannelThresholds[RCCL_AR_TUNABLE][0][0] = 1; // minByteThreshold (nonzero: 0 collides with CHAN_THRESHOLDS_UNDEFINED)

// minByteThreshold (nonzero: 0 collides with CHAN_THRESHOLDS_UNDEFINED)

comm->minMaxChannelThresholds[RCCL_AR_TUNABLE][0][0] = 1;

github-actions Bot added the project: rccl label Aug 28, 2026

aduforest marked this pull request as ready for review August 28, 2026 19:31

aduforest requested review from a team August 28, 2026 19:31

systems-assistant Bot added the organization: ROCm label Aug 28, 2026

aduforest force-pushed the users/aduforest/rccl-wrap-cc-tests-batch3 branch from 7c11694 to fe5d4e9 Compare August 28, 2026 23:50

aduforest and others added 12 commits August 28, 2026 17:15

aduforest force-pushed the users/aduforest/rccl-wrap-cc-tests-batch3 branch from fe5d4e9 to 21c0fb4 Compare August 29, 2026 00:33

amd-argus Bot reviewed Aug 30, 2026

View reviewed changes

Navigation Menu

Uh oh!

test(rccl): extend rccl_wrap.cc host coverage to 24 more functions (52.6% lines) - #10909

test(rccl): extend rccl_wrap.cc host coverage to 24 more functions (52.6% lines)#10909
aduforest wants to merge 12 commits into
users/aduforest/rccl-wrap-cc-tests-v2from
users/aduforest/rccl-wrap-cc-tests-batch3

aduforest commented Aug 28, 2026

Uh oh!

therock-pr-bot Bot commented Aug 28, 2026

Uh oh!

github-actions Bot commented Aug 29, 2026

Uh oh!

amd-argus Bot left a comment

Uh oh!

amd-argus Bot Aug 30, 2026

Uh oh!

amd-argus Bot Aug 30, 2026

Uh oh!

amd-argus Bot Aug 30, 2026

Uh oh!

amd-argus Bot Aug 30, 2026

Uh oh!

amd-argus Bot Aug 30, 2026

Uh oh!

amd-argus Bot Aug 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

	comm->minMaxChannelThresholds[RCCL_AR_TUNABLE][0][0] = 1; // minByteThreshold (nonzero: 0 collides with CHAN_THRESHOLDS_UNDEFINED)
	// minByteThreshold (nonzero: 0 collides with CHAN_THRESHOLDS_UNDEFINED)
	comm->minMaxChannelThresholds[RCCL_AR_TUNABLE][0][0] = 1;

Uh oh!

Conversation

aduforest commented Aug 28, 2026

What

How

Seams (process-isolation note)

Mutation

Test plan

Uh oh!

therock-pr-bot Bot commented Aug 28, 2026

✅ All Policy Checks Passed

Uh oh!

github-actions Bot commented Aug 29, 2026

RCCL Perf-Regression Gate: ⚠️ NO VERDICT (not measured)

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

amd-argus Bot Aug 30, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 30, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 30, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 30, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 30, 2026

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot Aug 30, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant