test(rccl): host microtests for four uncovered init.cc helpers - #10807
Open
pvallem wants to merge 2 commits into
Open
test(rccl): host microtests for four uncovered init.cc helpers#10807pvallem wants to merge 2 commits into
pvallem wants to merge 2 commits into
Conversation
JIRA ID : AICOMRCCL-1685 Covers four functions that were at zero line coverage, each taken to 100%: | function | tests | mutants killed | |---|---|---| | setCommAbortFlags (init.cc:3884) | 6 | 17/19 (2 documented equivalents) | | commCleanup (init.cc:3696) | 6 | 14/14 | | copyCommConfig (init.cc:3161) | 5 | 15/15 | | ncclGetUniqueId_impl (init.cc:358) | 9 | 17/17 | Mutation, not coverage, is the acceptance gate: each block was proven by mutating the code it covers and confirming the tests die. 63 of 65 mutants killed; the two survivors are argued equivalent in place (a store reordering between provably distinct locations, and a memory_order relaxation that needs a second thread to observe). New seams follow the existing catalogue and are reset in ResetInitFakes(): g_ncclCeFinalizeResult, g_ncclTunerPluginUnloadResult (+LastComm), g_cleanupCallOrder, g_ncclOsTopoGetStrFromSysResult (+Calls), g_ncclEnvPluginInitResult, g_bootstrapGetUniqueIdCalls, g_bootstrapHandleTemplate, and Recorder argument recording. Two fixes to fakes that were hiding coverage: - ncclParamGraphStreamOrdering() was hardcoded to 0, so envConfigOverride forced config.graphStreamOrdering = 0 on every call and the last field of ncclConfig_t was unobservable -- a memcpy short by 4 bytes would have passed. It now reads through the g_loadParam seam with production's default. - The bootstrapGetUniqueId fake wrote only `magic`, leaving the rest zero, so a shortened memcpy in ncclGetUniqueId_impl was invisible. It now writes a byte-wise distinctive handle. src/init.cc is untouched. Latent defects found while covering these blocks are pinned as current behaviour with a LATENT BUG note naming the line, not fixed: init.cc:3889 (childAbortFlagDev guarded by childAbortFlag's null check), init.cc:369 (Recorder::record return discarded where six sibling call sites NCCLCHECK it), init.cc:3698 (a cudaSetDevice failure aborts teardown before commFree, leaking the communicator; and the caller's current device is never restored), init.cc:3700 (a tuner plugin returning non-success leaks the comm). Verified: 283/283 and 282/282 uncached, rccl-UnitTestsMicro 30/30, rccl-HostUnitTests 282/282, clean under --gtest_shuffle --gtest_repeat=3 across three seeds, --gtest_death_test_style=threadsafe, and a hostile environment.
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
sdarisala
approved these changes
Aug 28, 2026
…elpers Conflict in test/host/init-test.cc: develop (#10808) lifted the inline NVTX redirector block into fakes/nvtx_redirect.h, and this branch added the MicroMalloc and tuner-finalize seams immediately above it. Kept both: the new seams stay, the inline block is replaced by the shared header include.
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 |
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-1685
What
Four
src/init.cchelpers were at zero line coverage. This takes each to 100%, adding 26 host-only microtests.src/init.ccis untouched — test-only.setCommAbortFlags():3884commCleanup():3696copyCommConfig():3161ncclGetUniqueId_impl():358init.ccfunctions covered: 57.97% → 62.32%.Mutation is the acceptance gate, not coverage
Every block was proven by mutating the code it covers and confirming the tests die — 63 of 65 mutants killed. Coverage alone does not show whether a suite can see a wrong answer; earlier work on this file repeatedly hit suites sitting at 100% line and branch coverage with live mutants.
The two survivors are argued equivalent in place rather than papered over: a store reordering between two provably distinct locations with no intervening observer, and a
memory_order_release→relaxedrelaxation that needs a second thread and TSan to observe.Two fakes were hiding coverage
ncclParamGraphStreamOrdering()was hardcodedreturn 0, soenvConfigOverrideforcedconfig.graphStreamOrdering = 0on every call and the last field ofncclConfig_twas unobservable — amemcpyshort by 4 bytes would have passed. Now reads through theg_loadParamseam with production's default.bootstrapGetUniqueIdfake wrote onlymagic, leaving the rest zero, so a shortenedmemcpyinncclGetUniqueId_implwas invisible. Now writes a byte-wise distinctive handle.New seams
All follow the existing catalogue and are reset in
ResetInitFakes(), with the reset verified load-bearing (deleting it makes tests fail under--gtest_shuffle):g_ncclCeFinalizeResult,g_ncclTunerPluginUnloadResult(+LastComm),g_cleanupCallOrder,g_ncclOsTopoGetStrFromSysResult(+Calls),g_ncclEnvPluginInitResult,g_bootstrapGetUniqueIdCalls,g_bootstrapHandleTemplate, andRecorderargument recording.commCleanup's oracle is a call-order vector, since ordering (finalize → unload → free) is the thing worth testing there and a return code cannot show it.Latent defects — pinned, not fixed
Found while covering these blocks. Current behaviour is pinned with a
LATENT BUGnote naming the line;src/is unchanged.childAbortFlagDevis dereferenced under a null check on a different pointer (childAbortFlag). Safe only because the two are assigned as a pair — but:2832and:4077both nullchildAbortFlagalone, andgroup.cc:584repeats the shape. Retargeting the guard to the more natural-looking pointer turnsncclCommAbortinto a write through a dangling device pointer.Recorder::instance().record(...)'s result is discarded, while all six sibling call sitesNCCLCHECKit. A recorder failure while minting a unique id is reported asncclSuccess, silently breaking replay for exactly the run being captured.cudaSetDevicefailure returns beforecommFree, leaking the whole communicator (bootstrap sockets, channels,sharedRes, proxy state). The only caller,commReclaim:3812, justWARNs and moves on.commCleanupnever restores the caller's device, soncclCommDestroysilently mutates the calling thread's current HIP device.mem_manager.cc:1094does save/restore for the same pattern.commFree— plugin code decides whether RCCL leaks a communicator.Test plan
rccl-UnitTestsMicroInitrccl-UnitTestsMicroInit-uncachedrccl-UnitTestsMicrorccl-HostUnitTests--gtest_shuffle --gtest_repeat=3, seeds 1/42/777--gtest_repeat=10--gtest_death_test_style=threadsafeNCCL_CTA_POLICY,NCCL_HOSTID,HSA_*,ROCSHMEM_HEAP_SIZE,NCCL_DEBUG,NCCL_TOPO_DUMP_FILE)No CMake or YAML change: all tests are in the existing
InitMicrotest/InitMicrotestIsolatedsuites, already wildcarded bytest_categories_micro_init.yaml. Built host-only against the hipified snapshot — no GPU, nolibrccl.so, no HIP runtime.