test(rccl): add host-only microtests for src/enqueue.cc - #10861
Open
ankith117 wants to merge 9 commits into
Open
test(rccl): add host-only microtests for src/enqueue.cc#10861ankith117 wants to merge 9 commits into
ankith117 wants to merge 9 commits into
Conversation
JIRA ID : AICOMRCCL-2198
src/enqueue.cc had zero host-side coverage: no ENQUEUE_CC_PATH macro, no test
referencing it, and none of its symbols in any host test binary. Coverage was 0%
by construction, not by measurement.
Adds rccl-UnitTestsMicroEnqueue with 171 tests covering the host-reachable
helpers. Test-only; src/ is untouched.
Coverage of enqueue.cc: 0% -> 24.06% lines, 42.59% functions, 12.39% branches.
15 functions reach 100%. Note the epic's 80% target is not achievable whole-file
for this unit: only 1028 of 2995 countable lines (34.3%) live in host-reachable
functions; the remainder is kernel-launch and orchestration (addP2pToPlan,
scheduleCollTasksToPlan, taskAppend, ncclPrepareTasks, ncclLaunchPrepare) that
needs a real communicator, HIP runtime and proxy threads. 24.06% whole-file is
roughly 70% of the reachable subset.
Mutation is the acceptance gate, not coverage. 45 mutants designed against the
covered blocks, 44 killed. The survivor is equivalent and argued in place:
dropping the count == 0 term from rcclKernelPackedChannels' early-out is masked
by the cellsPerChannel == 0 guard two lines later, which returns the same value
for the same input.
Mutation caught six tests that coverage could not. Four executed the mutated line
but could not observe its effect because another guard fired first: coll work
exceeding the 192-byte batch budget, the nNodes<=2 one-p2p-per-batch cap,
ncclDevWorkStorageTypeArgs being the zero value, and the 5120-byte kernel-args
floor. Two scaling tests used EXPECT_GE, which accepts no scaling at all.
Two structural notes:
* enqueue.cc:28 includes src/device/common.h, a device header that cannot
compile under --offload-host-only. Its include guard is pre-set and the six
ncclDevKernel_Generic_N kernels are supplied as host surrogates -- every use
is plan->kernelFn = ncclKerns[i].kernelFn, an opaque void* that is stored,
never dereferenced or launched. This covers host-side table indexing only;
kernel ABI and linkage are out of scope for a host-only binary, and the
limits are documented at the shim.
* Three shared nccl_stubs.cc entries are omitted for this target via
per-symbol RCCL_STUBS_OMIT_<symbol> macros, with the 1:1 mapping documented
in fakes/enqueue_stub_overrides.cc. enqueue.cc defines
ncclInitKernelsForDevice and ncclParamGraphStreamOrdering itself;
rcclUseAinic needs a real value rather than a fail-loud stub. The init and
p2p targets are unaffected and still pass.
Latent defects found and pinned, not fixed:
* :2660 rcclOverrideChannels is called with no NCCLCHECK, silently discarding
an ncclResult_t -- unlike every other result-returning call in
topoGetAlgoInfo. A failing channel override is ignored and tuning proceeds
on the unmodified value.
* :1550 calcP2pChannelCount is dead code: defined once, called nowhere.
* :2813 the NVLS CTA_POLICY_EFFICIENCY branch is unconditionally false
(recChannels = nMaxChannels + 1, then tested <= nMaxChannels), because
ncclNvlsRegResourcesQuery is commented out.
* :199 addProxyOpIfNeeded and :772 ncclAddProxyOpIfNeeded are near-identical
duplicate bodies that can drift; tests exercise the extern.
Investigated and cleared: the ncclAvg switch at :3232-3274 has no default arm,
which reads like an uninitialised-op hazard. It is unreachable -- ncclTypeSize
returns -1 for every type the switch would miss, so the nbits <= 0 guard at
:3210 rejects them first. A test pins that coupling so it fails loudly if either
side gains a datatype.
Known gaps, deliberate for this PR: no process-isolated suite for the four
function-local statics at :2528, :2529, :2686 and :2766 (each latches on first
call), no death tests on the divide-by-zero paths, and CaptureLog is used on 5
of the 12 INFO/WARN sites in the covered functions. Follow-up work.
Test plan:
rccl-UnitTestsMicroEnqueue 171/171
rccl-UnitTestsMicroInit 256/256
rccl-UnitTestsMicroInit-uncached 255/255
rccl-UnitTestsMicro (p2p) 30/30
--gtest_shuffle seeds 1/42/777, 3 repeats each: pass
ldd: no HIP, ROCm, HSA or RCCL linkage
Built host-only against the hipified snapshot -- no GPU, no librccl.so, no HIP
runtime. Reviewed with codex over three rounds; all findings addressed.
Co-Authored-By: Claude <noreply@anthropic.com>
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
JIRA ID : AICOMRCCL-2198 Follow-up to dee0fb1. Test-only; src/ remains untouched. 171 -> 176 tests. HIGH -- topoGetAlgoInfo was reading past a stack local Three tests in tests_batch9.inc passed &tablePtr, the address of an 8-byte loat* local, where topoGetAlgoInfo casts its loat** back to loat (*)[NCCL_NUM_PROTOCOLS] and reads all NCCL_NUM_ALGORITHMS * NCCL_NUM_PROTOCOLS cells -- 84 bytes past the end of it. Both production call sites pass (float**)collCostTable, the array itself. The initCollCostTable + ScriptAllTimes setup in those tests was therefore inert: the selected algorithm was decided by whatever stack bytes followed the local. It landed on TREE, which is the only algorithm that reaches rcclOverrideChannels, so the assertions passed by coincidence and any layout change would have flipped them. Fixed by reusing CostTable from tests_batch4.inc, whose ptr() already performs the correct cast. Added TopoGetAlgoInfo_SelectsTheOnlyScriptedCell, which scripts exactly one (algo, proto) cell and asserts selection follows it -- an assertion that cannot hold unless the table handed over is genuinely the one read. MEDIUM * tests_batch6.inc: BatchPlanComm constructed memScoped but not memPermanent, which ncclAddProxyOpIfNeeded allocates against. A value-initialised stack has bumper == end == 0, so allocate() fell through to allocateSpilled and malloc'd a 64 KiB hunk; because topFrame.hunk was null the hunk was never linked into the frame chain, so even a destruct could not reclaim it. Three tests took that path. Both stacks are now constructed and destructed. * tests_batch1.inc: the default-arm sweep for ncclFuncTrafficPerByte covered five hand-picked funcs and omitted ncclFuncAlltoAllvGda, so deleting case ncclFuncAlltoAllvGda: return nRanks; left every test green. Now sweeps the whole ncclFunc_t enum, which also stays closed when a func is added, plus a dedicated test for that arm. * tests_batch3.inc: the tree-depth ladder test used EXPECT_GE. Substituting 1 for tree.depth makes shallow and deep produce the same chunk size, so the mutant survived. Now EXPECT_GT (measured 131072 vs 65536). * tests_batch5.inc: g_rcclUseAinic had its own omit macro and its own override TU but no test ever wrote it, so dropping the !rcclUseAinic() conjunct -- the reason that file exists -- survived. Added the gfx950-with-AINIC case. * test_executor.py and ci-precheckin.json: rccl-UnitTestsMicroEnqueue was absent from the llvm-cov --object list and from both the test_configs and test_suites blocks. The micro binaries compile their unit under test in rather than linking librccl.so, so their counters live in the binary and cannot be attributed without an explicit --object entry. Without this the enqueue.cc coverage number never appears in a CI-generated report. The tests themselves already ran via the GitHub workflow. LOW * enqueue-test.cc: dropped the device_table.h include and the hipify/gensrc include dir it required. ncclDevKernelArgsDefaultStorage comes from src/include/device.h, already in scope via enqueue.h -> comm.h; the generated header supplied nothing this TU used. Verified by rebuilding both CMake modes. The shim comment now states where the typedef actually comes from. * CMakeLists.txt: the enqueue block now sets GENERATED TRUE on the three oracle TUs itself. It previously worked only because source-file properties are directory-scoped and the init block happens to run earlier in the file -- reordering the targets would have broken a fresh configure. * tests_batch8.inc: the unroll test compared two plans to each other, which rests on the six kernel surrogates having distinct addresses -- a property of the test shim. Now compares against ncclKerns[i].kernelFn, matching the sibling test and pinning the index arithmetic instead. * tests_batch1.inc: the gfx950 nthreads test is a tautology today because RCCL_GFX950_MAX_NTHREADS and RCCL_DEFAULT_MAX_NTHREADS are both 256. The mutant is genuinely equivalent, so the comment was the defect. Replaced with a static_assert that fails the build when the constants diverge, with a note to make the test differential at that point. * tests_batch5.inc: one SetParam(key, value) helper replaces SetBatchParam plus three inline re-rolls of the same lambda. * enqueue_fakes.h: about 23 seams were declared, documented and reset but never driven, which invites a reader to infer coverage that does not exist. Two carried comments promising a tested rejection path -- g_commEnsureReadyResult and g_recorderResult -- and those paths are now tested. The rest are marked LINK FLOOR ONLY. Also documented two hazards: the NCCL_PROTO latch at :2528 means --gtest_shuffle passing reflects a constant latch rather than proven order-independence and leaves the user-NCCL_PROTO arm unreachable here; and :2686 reads NCCL_PROTO/NCCL_ALGO through raw libc getenv, which g_enqEnv does not intercept. * enqueue-test.cc: documented that the nine tests_batchN.inc fragments share one anonymous namespace and that later ones use fixtures from earlier ones, so the include order is load-bearing, with a map of what each fragment holds. * MICROTEST_README.md: added the enqueue unit, which enqueue_fakes.h points readers at. NIT * Copyright headers added to the nine .inc files, matching their siblings. * enqueue_stubs.cc: removed an empty extern 'C' {} block and the FAIL_LOUD alias, calling Unreached() directly at all 14 sites. * test_categories_micro_enqueue.yaml: added comprehensive_regression and full_regression, which all seven existing category yamls carry. * .pre-commit-config.yaml: the clang-format hook's files regex did not cover .inc, so the bulk of this PR escaped format checking. * run_host_tests.sh and the CMakeLists NOTE block: stale binary lists updated. * enqueue-test.cc now states once that every enqueue.cc:NNNN citation refers to the hipified copy, which runs one line higher than src/enqueue.cc (verified: 4178 vs 4177 lines; rcclOverrideChannels at :2660 vs :2659). Test plan: rccl-UnitTestsMicroEnqueue 176/176 rccl-UnitTestsMicroInit 256/256 rccl-UnitTestsMicroInit-uncached 255/255 rccl-UnitTestsMicro (p2p) 30/30 --gtest_shuffle seeds 1/42/777, 3 repeats each: pass ldd: no HIP, ROCm, HSA or RCCL linkage Mutation re-run: 44 killed, 1 documented equivalent survivor Coverage of enqueue.cc: 24.06% -> 24.13% lines, 42.59% functions.
|
🎉 All checks passed! This PR is ready for review. |
…rd test Addresses the second-round review of the enqueue.cc host microtests. Merge tests_batch1..9.inc into enqueue-test.cc (3078 lines, alongside init-test.cc at 2979 and p2p-test.cc at 2814). The fragments were textual includes sharing one anonymous namespace with a load-bearing include order, which nothing else in the monorepo does. All 176 TEST_F cases are unchanged; the .inc files are removed. With no .inc files left, the clang-format regex in the root .pre-commit-config.yaml is dead config and is reverted, so the PR no longer touches anything outside projects/rccl. Fix TopoGetAlgoInfo_SelectsTheOnlyScriptedCell, which could not fail. It scripted g_topoGetAlgoTime and called topoGetAlgoInfo, but topoGetAlgoInfo never calls ncclTopoGetAlgoTime: the only call site is enqueue.cc:2539 inside updateCollCostTable. The fake was never invoked, every cost cell stayed at NCCL_ALGO_PROTO_IGNORE, and the two assertions only re-checked the RING/SIMPLE fallback installed at :2573-2580. Deleting the argmin loop body, inverting the comparison, or restoring the &tablePtr defect all left it green. The test now drives updateCollCostTable first and asserts TREE/LL, which is differential against that fallback. Renamed to SelectsTheCheapestScriptedCell. The three dead ScriptAllTimes calls in the sibling tests are removed. Also: - Restate the line-number banner in terms of src/enqueue.cc and renumber every enqueue.cc citation to that base. Twelve were already real-base while the banner told readers to subtract one, and seven were wrong in either base. - Replace the misplaced LINK FLOOR ONLY block with a per-declaration UNDRIVEN marker on each of the 21 undriven seams, so the claim travels with the declaration instead of a note that had drifted past the seams it named. - RedOpCreate_CommNotReady now uses ncclInternalError so it cannot pass on the ncclInvalidArgument that the preceding CommCheck also returns. - Correct the device.h include-chain comment, the fail-loud tier comment, the MICROTEST_README unit list and reset helper, and the CTest NOTE block.
Resolves the MICROTEST_README.md conflict in step 3 of "Adding a new test". develop rewrote that step to point at the new ScopedHook.h helper; this branch had generalised the reset-helper name because enqueue.cc is now a third unit under test. Both changes are wanted, so the resolution keeps develop's ScopedHook guidance and applies the generalised reset-helper list on top. develop's new host-test surface (rma-proxy-progress-test.cc, group-test.cc, comm_fakes, rma_fakes, collective_stubs) lands in TEST_MICRO_SOURCE_FILES, not TEST_MICRO_ENQUEUE_SOURCE_FILES, so the enqueue target picks up no new symbols and both of its source lists are unchanged by the merge.
Contributor
RCCL Perf-Regression Gate:
|
| 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 |
github-actions
Bot
added
the
github actions
pvallem
force-pushed
the
users/saverine/enqueue_microtest
branch
from
August 29, 2026 20:18
4486714 to
946db4a
Compare
rccl-UnitTestsMicroEnqueue fails to link in TheRock's build (clang 24) with
seven undefined symbols out of libclang_rt.profile-x86_64.a:
ld.lld: error: undefined symbol: __interception::DynamicLoaderAvailable()
ld.lld: error: undefined symbol: __sanitizer_internal_memset
ld.lld: error: undefined symbol: __prof_rocm::profRecordDrainedBounds(...)
>>> referenced by InstrProfilingPlatformROCm.cpp
The cause is not the coverage flags. rccl-UnitTestsMicro, MicroInit and
MicroInit-uncached all link with a byte-identical option set; only this target
fails. The discriminator is TU content:
1. enqueue.cc:2253 and :2339 call CUCHECKGOTO(cuLaunchKernel(...)). On AMD,
rocmwrap.h:96 defines CUPFN(symbol) as `symbol`, not `pfn_##symbol`, so
this is a direct call and hipify rewrites it to hipModuleLaunchKernel.
enqueue.cc names that symbol family 5 times; init.cc, p2p.cc and group.cc
name it zero times, which is exactly the set of targets that link.
2. Nothing defines it here: -no-hip-rt is on the link line and no file under
test/host/fakes/ supplies a HIP launch entry point.
3. -fprofile-instr-generate puts libclang_rt.profile on the line, and ROCm's
clang-24 fork adds InstrProfilingPlatformROCm.cpp.o to that archive, which
DEFINES hipModuleLaunchKernel (and 15 sibling HIP entry points) as
compiler-rt INTERCEPTORs.
4. lld extracts that member to satisfy the symbol, dragging in __interception
and __sanitizer_internal dependencies that ship in no archive on the line.
Member selection happens during symbol resolution, before --gc-sections,
so GC cannot prevent it.
It is invisible on ROCm 7.2.4 / clang 22 because that profile archive has no
such object, so nothing defines the symbol and --gc-sections then collects
ncclLaunchKernel's dead section without diagnosing the dangling reference.
Defining the stub keeps the target's stated contract -- an un-shimmed HIP call
must fail loudly rather than bind to whatever the linker happens to find, which
today is a profiling interceptor. No CMake change: enqueue_stubs.cc is already
in both source lists (test/host/CMakeLists.txt:252 and :730). Coverage flags
are untouched, so the llvm-cov wiring and the enqueue.cc coverage it reports
are unaffected.
enqueue-test.cc carried its own 18-line copy of the NVTX neutering block. develop has since extracted that exact block into fakes/nvtx_redirect.h (PR #10808) and both init-test.cc:54 and group-test.cc:66 now include it, so the develop merge left this branch holding the only remaining inline copy. The logic is identical; the header's comment is the generalised wording ("the unit-under-test's re-include" rather than "enqueue.cc's"). Include position is unchanged and still load-bearing: it must precede #include ENQUEUE_CC_PATH so nvtx.h is neutered before enqueue.cc re-includes it.
Generalises the hipModuleLaunchKernel stub to all 16 HIP entry points that ROCm's profile runtime also defines, and moves them into the shared fakes file. Background: with -fprofile-instr-generate, libclang_rt.profile-x86_64.a is on the link line, and ROCm's clang-24 fork adds InstrProfilingPlatformROCm.cpp.o to that archive. That object defines 16 HIP entry points as compiler-rt INTERCEPTORs. If a unit under test names one and nothing else defines it, lld extracts the member to resolve it and drags in __interception::*, __sanitizer_internal_* and __prof_rocm::* dependencies that ship in no archive on the line. Defining the symbol ourselves means the member is never consulted. Stubbing one symbol only moved the landmine: any future microtest whose unit names hipLaunchKernel, hipGraphLaunch, hipModuleLoad or one of the others hits the identical failure, and it is invisible outside TheRock CI. Local ROCm 7.2.4 / clang 22 ships no such archive member, so nothing defines the symbol and --gc-sections then drops the dead section without diagnosing the dangling reference -- which is exactly how this shipped. The list comes from compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp (llvm/llvm-project main, md5-identical on ROCm/llvm-project amd-staging; latest touching commit ae5e065a1128), cross-checked against installHipInterceptors(). It lives in fakes/hip_profile_interceptor_fakes.h as an X-macro list rather than as definitions: a header of definitions would either collide when included twice or, if marked inline, risk not being emitted when unreferenced -- which would silently reopen the bug. They go in hip_fakes.cc because that file's stated job is already "plain stubs for every other HIP symbol the object code references", it is disjoint from the 16, and it is already in every micro target's source list. So all four micro binaries are covered with no CMake change at all. Verified on a cleanly regenerated hipify tree (no borrowed tree): all four targets build, 16/16 symbols defined in each hip_fakes.cc.o, no duplicate symbols, and 176 + 110 + 259 tests pass. -Wl,-y confirms the definition is resolved from a regular object at link time. The archive-extraction half is unreproducible locally and rests on lld archive semantics plus upstream source.
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
There was a problem hiding this comment.
Automated Review Guard for Upfront Scrutiny
Findings
22 finding(s) are posted inline, on the lines they refer to.
projects/rccl/test/host/MICROTEST_README.md:448(Low) The standalone-build run block at MICROTEST_README.md:441-450 lists every micro binary except the new one
Bottom line: Nothing blocks this: the include shim, the RCCL_STUBS_OMIT_* contract and the CI wiring all check out. Worth fixing first is ShmemScratchWarpSize_NvlsTermOnlyAppliesAtArch900Plus at enqueue-test.cc:251, which cannot observe the gate it is named after.
Comment @argus review to re-run, or @argus re-review after pushing.
| ${PROJECT_BINARY_DIR}/hipify/src/misc/utils.cc | ||
| PROPERTIES GENERATED TRUE) | ||
|
|
||
| add_executable(rccl-UnitTestsMicroEnqueue ${TEST_MICRO_ENQUEUE_SOURCE_FILES}) |
There was a problem hiding this comment.
This is the third copy of the same micro-test target block in the if(BUILD_TESTS) branch: the foreach(_hip_tgt hip::host hip::device) block is verbatim at 53-62, 187-196 and 293-302, the BUILD_ADDRESS_SANITIZER block at 97-103, 209-215 and 312-318, and apply_test_category_labels at 118-131, 217-230 and 320-334. The standalone branch repeats the same shape at 609-653, 680-720 and 744-781. Would a rccl_add_micro_test(NAME ... SOURCES ... DEFS ... CATEGORIES_YAML ...) function collapse them? That removes roughly 60 lines here and 30 in the standalone branch, and makes the next microtest target one call rather than a fourth paste.
| } | ||
|
|
||
| // =========================================================================== | ||
| // addProxyOpIfNeeded (enqueue.cc:198) / ncclAddProxyOpIfNeeded (:771) |
There was a problem hiding this comment.
This header names addProxyOpIfNeeded (enqueue.cc:199) and ncclAddProxyOpIfNeeded (enqueue.cc:772), but all five tests below call only the extern. The two bodies are byte-identical and the static is the one the launch path uses (enqueue.cc:882 and :1053), which is exactly the drift risk the PR description raises. Since the unit is #included, the static is directly callable here. Could one differential test drive both with the same ncclProxyOp and assert the same queue state, so a divergence fails instead of going unnoticed?
| // This pins today's behaviour (src/ is untouched): the failure does NOT | ||
| // propagate. If a future change adds the missing NCCLCHECK, this test fails and | ||
| // should become an assertion that the error IS returned. | ||
| TEST_F(EnqueueMicrotest, TopoGetAlgoInfo_OverrideChannelsFailure_IsSilentlyIgnored) { |
There was a problem hiding this comment.
The PR description lists enqueue.cc:2813, the NVLS CTA_POLICY_EFFICIENCY branch that is unconditionally false because recChannels = nMaxChannels + 1 is then tested <= nMaxChannels, among the four defects pinned with a test. Grep here for CTA_POLICY, recChannels, RegResourcesQuery and EFFICIENCY finds nothing; the other three pins are all present and labelled. Could the description drop that row, or a one-line pin be added next to this one?
| TEST_F(EnqueueMicrotest, RedOpCreate_FirstCall_GrowsCapacityToFour) { | ||
| // `if (cap < 4) cap = 4` -- the first allocation jumps straight to 4, it does | ||
| // not double from 0. | ||
| RedOpComm rc; |
There was a problem hiding this comment.
The same four-line RedOpComm rc; create preamble appears ten times (1625, 1637, 1654, 1667, 1749, 1764, 1784, 3039, 3053) and the int(ncclUserRedOpMangle(rc.get(), op)) - int(ncclNumOps) slot arithmetic six (1642, 1659, 1672, 1701, 1718, 1726). The scalar local is scalar in four tests and s in nine, for the same thing. Would ncclRedOp_t CreateHostImmediate(RedOpComm&, float) plus int SlotIndex(RedOpComm&, ncclRedOp_t) cover both? About 30 lines, and the mangle arithmetic then lives in one place.
| // basename-unique in the tree, so hipify keeps its name (no _tmp suffix). | ||
| #include ENQUEUE_CC_PATH | ||
|
|
||
| class EnqueueMicrotest : public ::testing::Test { |
There was a problem hiding this comment.
The PR description reports 171 tests and a 171/171 test-plan result, but this file defines 176 TEST_F, all in suite EnqueueMicrotest. The only preprocessor guard is #ifndef RCCL_DEVICE_LINKER at 254-281 covering three of them, and that macro is set only at src/CMakeLists.txt:1065 and test/CMakeLists.txt:800, neither of which reaches this target, so all 176 build. Could the description be refreshed so the reported pass count matches the binary?
Ten fixes from the automated review. Every factual claim checked out, including
two against code added earlier in this branch.
Defects:
- hip_profile_interceptor_fakes.h and enqueue_stubs.cc both pointed at
hip_profile_interceptor_fakes.cc, which does not exist; the interceptor
definitions live in hip_fakes.cc. That comment is the map for whoever hits
the next __interception::* link error, so a wrong pointer is expensive.
- enqueue_stubs.cc carried #include <hip/hip_runtime.h>, added only for the
hipModuleLaunchKernel stub that has since moved to hip_fakes.cc. Deleted
rather than reordered; the file compiled without it before and does again.
- CalcCollChunking_Reduce_TreeVsRing dropped the two ASSERT_EQ(ncclSuccess)
lines its Broadcast twin has, so it read proxyOp.pattern unchecked. A
calcCollChunking error for Reduce would have surfaced as a pattern mismatch
rather than as the error it is.
- Three RCCL_STUBS_OMIT_ comments in nccl_stubs.cc exceeded the 120-column
limit (127/126/139). Wrapped; the three remaining long lines predate this PR.
Claims that were not true:
- The ncclKerns[] shim rationale said those addresses are never dereferenced
or launched. enqueue.cc:110 feeds them to cudaFuncGetAttributes (:115) and
cudaFuncSetAttribute (:121), and :2218 reads plan->kernelFn to launch. The
true and sufficient claim is that no covered path reaches them.
- The rcclKernelPackedChannels comment presented the `count == 0` term as
covered. It is an equivalent mutant: at count == 0, cells is 0, so
cellsPerChannel is 0 and the :190 guard returns the same nMaxChannels.
- enqueue_fakes.h claimed the min/max channel defaults mirror production.
Production is -2 for both (graph/connect.cc:831-832); the fakes hold
already-resolved clamps. Comment corrected, values deliberately left alone.
Names that promised more than the assertion delivers:
- ShmemScratchWarpSize_NvlsTermOnlyAppliesAtArch900Plus was an equivalent
mutant -- ncclNvlsUnrollBytes is 4*16 for every arch and ncclCollUnroll is
never below 4, so SIMPLE ties or wins everywhere. Renamed to
_NvlsTermNeverWins and backed with a static_assert, matching the treatment
the Gfx950 test beside it already had.
- PackedChannels_LLProtocolQuadruplesTraffic asserts only EXPECT_GE, which an
implementation dropping trafficPerByte *= 4 satisfies with equality. The
strict check already lives in the _Isolated test, so this is a rename to
_LLNeverNeedsFewerChannelsThanSimple.
- Ll128GrainDependsOnWarpSizeAndLineElems never varies either input.
Coverage:
- AddWorkBatch_P2pIneligibleAfterEligible passed batchP2P=false on both calls,
so the `newBatch |= !chan->wipBatch.batchP2P` conjunct at enqueue.cc:227 was
never the splitter -- the nP2ps == 1 else-arm split on its own. Verified by
running the mutation: with (false, false) the test passed with the conjunct
deleted; with (false, true) it fails. Renamed to _P2pEligibleAfterIneligible.
Not addressed here, deliberately: the review's consolidation suggestions (a
rccl_add_micro_test CMake function, deduplicating the comm fixtures, TEST_P over
the pattern tables, extracting call-site helpers). They are fair, but they are
~200 lines of untestable restructuring of someone else's test file, and a
refactor that green tests cannot validate is exactly how the &tablePtr defect
reached this branch. Better as a follow-up whose diff reads as a refactor.
Verified on a regenerated hipify tree: 176 + 110 + 259 tests pass, src/ untouched.
| finishPlan(f.c(), f.p()); | ||
| ASSERT_NE(nullptr, f.p()->kernelArgs); | ||
| EXPECT_EQ(f.p()->channelMask.masks[0], f.p()->kernelArgs->channelMask.masks[0]); | ||
| EXPECT_EQ(f.p()->workStorageType, f.p()->kernelArgs->workStorageType); |
There was a problem hiding this comment.
The workStorageType half of this assertion is 0 == 0 by construction, so deleting plan->kernelArgs->workStorageType = plan->workStorageType at enqueue.cc:316 kills no test. Production sets workStorageType to ncclDevWorkStorageTypeArgs, enumerator 0 (device.h:650), and kernelArgs comes from ncclMemoryStackAlloc (enqueue.cc:313), which memsets the block to 0 (utils.h:283). Nothing else in this file reads kernelArgs->workStorageType. Could FinishPlan_OversizedPlan_KeepsFifoStorage at 2334 carry it instead, where the stamped value is non-zero? The channelMask half on 2368 is sound.
| // `newBatch |= batch->workType != workType`. | ||
| BatchPlanComm bp; | ||
| addWorkBatchToPlan(bp.c(), bp.p(), 0, ncclDevWorkTypeColl, 7, 0); | ||
| addWorkBatchToPlan(bp.c(), bp.p(), 0, ncclDevWorkTypeCollReg, 7, 0); |
There was a problem hiding this comment.
AddWorkBatch_DifferentWorkType_ForcesNewBatch is split by the byte budget, not by the workType guard at enqueue.cc:220 it is named after. Mutating that line to newBatch |= false leaves the test green:
sizeof(ncclDevWorkColl) = 160 B
NCCL_MAX_DEV_WORK_BATCH_BYTES = 192 B (device.h:556)
after the Coll call, workBytes = 160 B
enqueue.cc:242 192 < 160 + sizeof(ncclDevWorkCollReg) -> new batch regardless
A p2p item followed by a Bcast reaches the nBcasts == maxitem arm at enqueue.cc:238 instead of the byte budget, leaving enqueue.cc:220 as the only splitter. Could the second call be ncclDevWorkTypeBcast after a p2p, keeping the offset a multiple of ncclDevWorkSize(ncclDevWorkTypeBcast) so enqueue.cc:246 does not open an extension batch?
| ASSERT_EQ(ncclSuccess, updateCollCostTable(cc.get(), &task, 1 << 20, /*collNet=*/0, | ||
| /*nvls=*/0, 1, 0, tbl.ptr())); | ||
| for (int p = 0; p < NCCL_NUM_PROTOCOLS; ++p) { | ||
| EXPECT_FALSE(tbl.written(NCCL_ALGO_NVLS, p)) << "p=" << p; |
There was a problem hiding this comment.
UpdateCollCostTable_NvlsUnsupported_SkipsNvlsAlgorithms asserts only the negative, so making the NVLS skip at enqueue.cc:2503-2505 unconditional kills no test. The two call sites that pass nvlsSupport=1 (1428 and 1444) take the short circuit at enqueue.cc:2491 and return before the algorithm loop, which those tests confirm themselves with EXPECT_EQ(0, g_topoGetAlgoTimeCalls). Its CollNet sibling UpdateCollCostTable_TooManyLocalRanks_SkipsCollNetAlgorithms at 1466 pairs the negative with a control at 1483-1489 that flips the input and asserts the row IS populated. Could the same control be added here?
| TEST_F(EnqueueMicrotest, GetImplicitOrder_CapturingIsIrrelevantOnAmd) { | ||
| // The `capturing` parameter only matters inside the #if'd-out CUDA block, so | ||
| // on AMD both values must give the same answer. Pins that the AMD arm really | ||
| // is unconditional -- if the #if ever changes, this fails. |
There was a problem hiding this comment.
GetImplicitOrder_CapturingIsIrrelevantOnAmd cannot fail if the #if at enqueue.cc:1998 changes, so this line overstates it. Under the seam value ncclCudaDriverVersionCache = 12000 (fakes/enqueue_fakes.cc:315, read by ncclCudaDriverVersion at rocmwrap.h:217) the CUDA arm also returns Serial for both values:
driver = 12000
capturing=true enqueue.cc:2002 12000 < 12090 -> Serial
capturing=false enqueue.cc:2006 12030 <= min(CUDART, 12000) -> false -> Serial
getImplicitOrder's third parameter int driver = -1 (enqueue.cc:1996) is the input that separates the arms, and no test here passes it. Could a case pass an explicit driver, or the claim be narrowed to what the assertions cover?
| op.opCount = 0x77; | ||
| g_proxySaveOpJustInquire = false; | ||
| ASSERT_EQ(ncclSuccess, ncclAddProxyOpIfNeeded(bp.c(), bp.p(), &op)); | ||
| EXPECT_TRUE(g_proxySaveOpSawJustInquireIn) |
There was a problem hiding this comment.
g_proxySaveOpSawJustInquireIn pins the fake rather than a production contract. Real ncclProxySaveOp runs if (justInquire) *justInquire = false; as its first statement (proxy.cc:631), so the bool needed = true seed at enqueue.cc:199 and enqueue.cc:772 is never read and flipping it to false changes nothing in production while failing here. What signals inquiry is passing a non-null pointer, not the pointee. Could this assertion go, leaving the channel check on 2748?
| // =========================================================================== | ||
|
|
||
| TEST_F(EnqueueMicrotest, ShmemScratchWarpSize_IsSixteenByteAligned) { | ||
| // The `+15 & -16` pad is the last operation; dropping it changes these. |
There was a problem hiding this comment.
The +15 & -16 pad at enqueue.cc:67-68 is an equivalent mutant, so this line overstates what the test catches. Every term entering the max is already a multiple of 16:
LL 0
LL128 64 * WarpSize (NCCL_LL128_SHMEM_ELEMS_PER_THREAD = 8, device.h:203)
SIMPLE (ncclCollUnroll * WarpSize + 1) * 16
NVLS 64 * WarpSize + 16 (ncclNvlsUnrollBytes = 4*16, device.h:741)
ShmemScratchWarpSize_NvlsTermNeverWins 25 lines below already gets the HONEST SCOPE plus static_assert treatment for exactly this shape. FinishPlan_SmallPlan_UsesArgsStorage at 2325 is the one other unannotated case, though 2831-2836 records it at the reinforcement site. Could both get the same note?
| TEST_F(EnqueueMicrotest, InitCollCostTable_MarksEveryCellIgnored) { | ||
| float table[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS]; | ||
| // Poison every cell first: a partial fill must be detectable. | ||
| for (auto& row : table) for (auto& v : row) v = -1.0f; |
There was a problem hiding this comment.
Unbraced outer for nesting an inner for. docs/dev_guide/nccl_coding_style.md:170 requires braces on all but the inner-most nested control statement, and the guide's not-OK example at 193-195 is this shape. Same line at 1339.
Suggested change
| for (auto& row : table) for (auto& v : row) v = -1.0f; | |
| for (auto& row : table) { | |
| for (auto& v : row) v = -1.0f; | |
| } |
| AvgComm comm(4); | ||
| const struct { ncclDataType_t dt; bool sgn; } kCases[] = { | ||
| {ncclInt8, true}, {ncclInt32, true}, {ncclInt64, true}, | ||
| {ncclUint8, false},{ncclUint32, false},{ncclUint64, false}}; |
There was a problem hiding this comment.
Suggested change
| {ncclUint8, false},{ncclUint32, false},{ncclUint64, false}}; | |
| {ncclUint8, false}, {ncclUint32, false}, {ncclUint64, false}}; |
|
|
||
| TEST_F(EnqueueMicrotest, InitCollCostTable_DoesNotWritePastTheTable) { | ||
| // Guard rows on both sides catch a wrong stride in the (float(*)[N]) cast. | ||
| struct { float before[NCCL_NUM_PROTOCOLS]; |
There was a problem hiding this comment.
Suggested change
| struct { float before[NCCL_NUM_PROTOCOLS]; | |
| struct { | |
| float before[NCCL_NUM_PROTOCOLS]; | |
| float table[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS]; | |
| float after[NCCL_NUM_PROTOCOLS]; | |
| } buf; |
|
|
||
| EXPECT_EQ(1, g_rcclUpdateCollectiveProtocolCalls) << "rcclUpdateCollectiveProtocol (:2608)"; | ||
| EXPECT_EQ(1, g_rcclSetPipeliningCalls) << "rcclSetPipelining (:2610)"; | ||
| EXPECT_EQ(1, g_rcclUpdateThreadThresholdCalls) << "rcclUpdateThreadThreshold"; |
There was a problem hiding this comment.
Suggested change
| EXPECT_EQ(1, g_rcclUpdateThreadThresholdCalls) << "rcclUpdateThreadThreshold"; | |
| EXPECT_EQ(1, g_rcclUpdateThreadThresholdCalls) << "rcclUpdateThreadThreshold (:2649)"; | |
| EXPECT_EQ(1, g_rcclOptThreadBlockSizeCalls) << "rcclOptThreadBlockSize (:2731)"; |
|
Summary for the 10 inline comments just posted. It is a separate comment because the review was submitted with an empty body by mistake and GitHub does not allow adding one afterwards. Automated Review Guard for Upfront Scrutiny Findings12 finding(s) are posted inline, on the lines they refer to.
Fixed since the last review
Still open
New this round
Bottom line: Nothing blocks this, and 11 of the 12 items you marked fixed check out at head. The substance this round is four tests that cannot fail on the mutation they are named for, sharpest at enqueue-test.cc:2369, which asserts 0 == 0. Comment |
Comment on lines
+315
to
+318
| int ncclCudaDriverVersionCache = 12000; | ||
| bool ncclCudaLaunchBlocking = false; | ||
| int ncclProfilerEventMask = 0; | ||
| std::unordered_map<uint64_t, int> ncclDevFuncNameToId; |
There was a problem hiding this comment.
Suggested change
| int ncclCudaDriverVersionCache = 12000; | |
| bool ncclCudaLaunchBlocking = false; | |
| int ncclProfilerEventMask = 0; | |
| std::unordered_map<uint64_t, int> ncclDevFuncNameToId; | |
| int ncclCudaDriverVersionCache = 12000; | |
| bool ncclCudaLaunchBlocking = false; | |
| int ncclProfilerEventMask = 0; | |
| std::unordered_map<uint64_t, int> ncclDevFuncNameToId; | |
| // TODO: reset these four in ResetEnqueueFakes() alongside the seams. They are process | |
| // state enqueue.cc reads at enqueue.cc:2163 and enqueue.cc:3544, and nothing restores | |
| // them, so with --gtest_shuffle this is a seed-dependent flake. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : AICOMRCCL-2198
What
Host-only microtests for
src/enqueue.cc, which had zero host-side coverage: noENQUEUE_CC_PATHmacro, no test referencing it, and none of its symbols in any host testbinary. Its coverage was 0% by construction, not by measurement.
171 tests in a new
rccl-UnitTestsMicroEnqueuebinary. Test-only —src/enqueue.ccisuntouched.
finishPlan:295ncclRedOpCreatePreMulSum_impl:4104ncclGetCollNetSupport:2457rcclKernelPackedChannels:177ncclAddProxyOpIfNeeded:772initCollCostTable:2477geteActivationMask/gettaskEventHandle:1820getImplicitOrder:1997calcP2pChannelCount:1550rcclEffectiveP2pBatchEnable:1170ncclTestBudget:406rcclShmemScratchWarpSize/DynamicSize:60ncclPlanSetDefaultKernel:2867addWorkBatchToPlan:209ncclRedOpDestroy_impl:4148hostToDevRedOp:3186ncclFuncTrafficPerByte:158updateCollCostTable:2487calcCollChunking:2871waitWorkFifoAvailable:1652Whole-file
enqueue.cc: 0% → 24.06% lines, 0% → 42.59% functions, 12.39% branches.15 functions at 100%.
On the 80% target
The epic's acceptance criterion is 80% host coverage. That is not achievable whole-file for
this unit. Of 2995 countable lines, only 1028 (34.3%) live in host-reachable functions;
the remaining 65.7% is kernel-launch and orchestration —
addP2pToPlan(281 lines),scheduleCollTasksToPlan(273),taskAppend(186),ncclPrepareTasks(183),ncclLaunchPrepare(142) — which need a real communicator, the HIP runtime and proxy threads.24.06% whole-file is roughly 70% of the reachable subset. Reporting both denominators so
the number is not misread; the ceiling is structural, not a shortfall in effort.
How
Follows the pattern from #10182 / #10590 / #10807 (init.cc) and #10808 (devcomm): the test TU
#includes the hipifiedenqueue.ccsostatichelpers are reachable, and links neitherlibrccl.sonor the HIP runtime. Every external symbol comes fromfakes/; deep paths getfail-loud stubs.
One structurally new problem.
enqueue.cc:28includessrc/device/common.h, a deviceheader that cannot compile under
--offload-host-only(extern __shared__variables,undeclared device intrinsics). Neither
init.ccnorp2p.ccreferences a__global__symbol,so this was new ground for the harness.
Resolved by pre-setting that header's own include guard and supplying the six
ncclDevKernel_Generic_Nkernels as ordinary host functions. Safe because every use of thetable is
plan->kernelFn = ncclKerns[i].kernelFn— an opaquevoid*that is stored, neverdereferenced or launched on a host path. Pre-setting another header's guard is the same
technique
init-test.ccalready uses for an NVTX collision.The limits are documented at the shim: these are host surrogates, so kernel ABI, symbols
and device linkage are out of scope for a host-only binary. The table tests are named and
scoped accordingly.
Three shared stubs are omitted for this target, each via its own
RCCL_STUBS_OMIT_<symbol>macro rather than one target-wide switch, with the 1:1 mappingdocumented in
fakes/enqueue_stub_overrides.cc.enqueue.ccdefinesncclInitKernelsForDeviceandncclParamGraphStreamOrderingitself;rcclUseAinicneeds areal value rather than a fail-loud stub. The init and p2p targets are unaffected and still
pass.
Mutation is the acceptance gate, not coverage
45 mutants designed against the covered blocks, 44 killed. The survivor is argued
equivalent in place rather than papered over: dropping the
count == 0term fromrcclKernelPackedChannels' early-out is masked by thecellsPerChannel == 0guard two lineslater, which returns the same value for the same input.
Mutation caught six tests that coverage could not. Four executed the mutated line but could
not observe its effect, because a different guard fired first:
newBatch |= funcId != devFuncId→falsenewBatch |= p2pRound == p2pRounds[i]→falseworkStorageType = ArgsncclDevWorkStorageTypeArgsis the zero value, so a zero-init plan already read as ArgsalignUp(kernelArgsSize, 16)→8sizeof(ncclDevKernelArgsDefaultStorage)is 5120, already 16-aligned, and floors every small planTwo more asserted
EXPECT_GEon tests named "scales with", which an implementation ignoringnRanksor the element size entirely would still pass. All six are fixed and the correspondingmutants now die.
Latent defects — pinned, not fixed
Found while building this.
src/is unchanged; each is pinned with a test documenting currentbehaviour so a future change fails loudly.
rcclOverrideChannelsis called with noNCCLCHECK, silently discarding anncclResult_t— unlike every other result-returning call intopoGetAlgoInfo. A failing channel override is ignored and tuning proceeds on the unmodified value.calcP2pChannelCountis dead code — defined once, called nowhere in the repo.CTA_POLICY_EFFICIENCYbranch is unconditionally false:recChannels = nMaxChannels + 1, then tested<= nMaxChannels, becausencclNvlsRegResourcesQueryis commented out.addProxyOpIfNeededandncclAddProxyOpIfNeededare near-identical duplicate bodies that can drift. Tests exercise the extern.Investigated and cleared: the
ncclAvgswitch at :3232-3274 has nodefault:arm, whichreads like an uninitialised-
ophazard. It is unreachable —ncclTypeSizereturns-1forevery type the switch would miss, so the
nbits <= 0guard at :3210 rejects them first. A testpins that coupling so it fails loudly if either side gains a datatype.
Known gaps
Deliberate for this PR, called out rather than left to be discovered:
enqueue.cchas four function-local statics (:2528, :2529,:2686, :2766) that latch on first call; the alternate value is unobservable in-process. The
init tests use
RUN_ISOLATED_TESTfor exactly this.chunkSize / grainSizesites).CaptureLogis used on 5 of the 12INFO/WARNsites in the covered functions.Test plan
rccl-UnitTestsMicroEnqueuerccl-UnitTestsMicroInitrccl-UnitTestsMicroInit-uncachedrccl-UnitTestsMicro(p2p)--gtest_shuffle, seeds 1/42/777, ×3 repeatslddThe init/p2p regression runs matter because
fakes/nccl_stubs.ccis shared. Built host-onlyagainst the hipified snapshot — no GPU, no
librccl.so, no HIP runtime. Whole suite runs in~50 ms.