Repository navigation
test(rccl): cover every gin_host.cc entry point with host-only tests - #13225
Merged
Merged
Conversation
…-only tests PR #12475 added the first host-only coverage of `src/gin/gin_host.cc` (the GIN proxy NUMA affinity pin) and PR #12245 the GIN_PROXY_NTHREADS progress-thread behaviour, which together left the file at 59% of lines and 34% of branches. This builds out the rest, so every function in the file now has a test. The test TU's `ncclGin_t` fake grows a uniform `ScriptedFailure` per plugin entry point ("fail call N of this function"), recorders for the `ginConfig`s and `regMrSym`/`deregMrSym` arguments, and knobs for a NULL-returning `createContext` and a `queryLastError` that reports an error on a chosen call. That makes each failure arm reachable without a bespoke fake per test. The same mechanism now also covers `ncclTopoGetLocalGinDevs` and `bootstrapAllGather`, and the rail team's stride is a knob so the stride-validation branches are reachable. New suites, one per entry point: `GinHostTypeQueryMicrotest`, `GinHostSignalRequestMicrotest`, `GinHostConnectOnceMicrotest`, `GinHostBackendSelectMicrotest`, `GinHostDevCommSetupMicrotest`, `GinHostDevCommFreeMicrotest`, `GinHostFinalizeMicrotest`, `GinHostRegisterMicrotest` and `GinHostQueryLastErrorMicrotest`. They all start with `GinHost`, so the existing `GinHost*` CTest pattern picks them up without touching the categories yaml beyond its description. Coverage of `gin_host.cc` in `rccl-UnitTestsMicroGinHost` goes from 59.01% to 100% of lines, 66.67% to 100% of functions, and 33.79% to 74.43% of branches; no source-level branch is left uncovered, the remainder being `NCCLCHECK`/`WARN` macro expansions at call sites whose failure arm no test drives. Testing: `run_host_tests.sh all` passes (80 GinHost tests, shuffled); ten mutations of the hipified `gin_host.cc` -- dropping the NULL-window check, the multi-segment DMABUF requirement, the stride-multiple check, the empty-slot skip, the early return on a reported error, the context round-up, the connectivity check in `ncclGetGinType`, the VA-signal requirement check, the devComm reset on failure, and the fail-path `closeListen` -- each turn exactly the matching test red. Co-authored-by: Pi Agent
…kes files
Review item 1 on gin-host-test.cc: the suite carried a second, weaker GIN
backend double alongside FakeGin, and defined four externals locally that
other fakes/ files already own.
Fold GinHostProxyAffinitySetupMicrotest into GinHostDevCommSetupMicrotest.
Its single test drove ginDevCommSetupWithBackend through a bespoke
FakeGinBackend vtable and rebuilt the ncclComm / ncclSharedResources /
backend state GinHostTest::SetUp already builds. On FakeGin it is three
lines, so FakeGinBackend, FakeCreateContext and FakeDestroyContext are gone.
Move the externals to the production TU that defines each, per
MICROTEST_README.md "Where a fake belongs":
bootstrapAllGather -> fakes/bootstrap_stubs.cc (seam already
existed; the fixture installs the
connection-count exchange on it)
ncclOsCpuCount / ncclOsSetAffinity -> fakes/os_fakes.cc (which already had a
parallel pair of recorders)
ncclSetThreadName -> fakes/nccl_fakes.cc, alongside the other
src/debug.cc symbols; drops the copy in
ras-test.cc
ncclTeamWorld / ncclTeamRail / -> fakes/nccl_device_core_fakes.cc, a new
ncclTeamRankToWorld file named for src/nccl_device/core.cc
The team accessors were in dev_runtime_micro_fakes.cc, named for a TU that
does not define them; the new file takes ncclTeamCft and ncclTeamCftMultimem
with them. ncclTeamLsa stays in devcomm_fakes.cc because its default silently
satisfies the comparison devcomm-test.cc's filters make and the trap comment
saying so is load-bearing; that exception is now recorded in the README.
os_fakes.cc's affinity recorders are reached from spawned proxy progress
threads, so the call counter becomes atomic and the two mask vectors are
written under a mutex.
For signature drift: the moved seams get ASSERT_HOOK_MATCHES_PROD from their
new homes. The nccl_device accessors are each overloaded with a device entry
point, so `&ncclTeamWorld` is ambiguous; ASSERT_HOOK_MATCHES_SIG in
signature-drift.h names the host overload instead, replacing the hand-rolled
static_assert devcomm_fakes.cc carried for ncclTeamLsa. ncclParamGinType and
ncclTopoGetLocalGinDevs stay in the test TU (neither owner's fakes file can
link into this binary), where the visible production declarations already make
a signature change a compile or link error -- verified by mutating each.
rccl-UnitTestsMicroGinHost: 80 tests, 155 ms, gin_host.cc still 100% line and
function coverage. All twelve host-only binaries green.
Co-authored-by: Pi Agent
…hared headers
Four helpers in gin-host-test.cc were generic test-side scaffolding that other
microtests hand-roll (or will need): the "fail call N of this entry point"
script, the bounded wait for another thread, the worker-thread joiner, and a
scope guard for cleanup that must survive a fatal ASSERT. Move each into its
own header next to ScopedHook.h, following the same one-helper-per-header
pattern:
ScriptedFailure.h ScriptedFailure, unchanged behaviour.
WaitUntil.h waitUntil, with the 2s budget as a named constant.
JoinThreadsOnExit.h the former JoinProgressThreads, with the GIN-specific
stopProgress() turned into a stop callback so it can
join any set of workers. gin-host-test binds the
callback in a progressWorkers() factory.
ScopeExit.h replaces the hand-written FinalizeGuard.
group-test.cc's SpinUntilAbortFlagObserved now uses waitUntil instead of its
own deadline loop.
Also dedup the four copies of the "finalize memsets the state, restore its C++
lifetime" placement-new into a GinHostTest::restoreGinStateAfterFinalize()
helper that states the reason once.
MICROTEST_README gains a "Shared test-side helpers" section listing the
headers and when to reach for each, so the next test promotes rather than
copies.
rccl-UnitTestsMicroGinHost: 80 tests, 173 ms, all passing.
rccl-UnitTestsMicro: 1655 tests, all passing.
Co-authored-by: Pi Agent
…left open Review item 3: three groups of gin-host-test.cc tests asserted little beyond the fake they were driving. The *Propagates error arms only checked that the status they injected came back. Each now also pins what must not happen on that branch, with the output pre-poisoned where "not written" would otherwise be indistinguishable from a zero-initialised field: devices/topology failures leave ginCommCount and the plugin alone, the connection-count exchange opens nothing, a failed closeColl leaves the connection slot and the GIN state intact, and a failed register/deregister/queryLastError stops at the failing connection instead of walking on. ncclGinQueryLastError's up-front clear of *hasError was being supplied by the fake, which also wrote the out-parameter on its failure path; the fake now reports nothing when it fails, so the assertion is on production. ExclusiveContextsAreConfiguredLikeSharedOnes pinned a non-behaviour -- it asserted the shared-path numbers, so it passed whatever the flag did. It is now ExclusiveContextsAreNotSharedWithAnotherDevComm, which asserts what the flag means: a second request gets its own contexts. The ncclGinDevCommFree and ncclGinQueryLastError suites walked a ncclGinStateDevComm list the test had calloc'd itself, so they partly tested the helper's model of that list. Both now build their devComms through ncclGinDevCommSetup. The hand-built list survives only for the ncclGinProgress tests, which need a different needsProxyProgress per connection, and appendDevComm folds into its only caller. rccl-UnitTestsMicroGinHost: 80 tests, 155 ms, gin_host.cc coverage unchanged (83.72% region / 74.43% branch). Each strengthened assertion was checked red by mutating the hipified gin_host.cc. Co-authored-by: Pi Agent
…h the public entry GinHostDevCommSetupMicrotest reached ginDevCommSetupWithBackend directly for every case, justified by ncclGinDevCommSetup collapsing each per-backend failure into ncclInternalError. That only holds for the refusals that have to be seen as a specific status; the tests that succeed, or that return ncclInternalError anyway, observe createdConfigs / devComm_ fields that are just as visible through the public call with the single backend the fixture configures -- and going through it exercises the type and signal gating in front of the helper for free. Ten tests now call ncclGinDevCommSetup: the backend-version table, context-count rounding, exclusive contexts, traffic class, signal pass-through, both stride cases, progress-thread reuse, the affinity stash, and the incomplete-context guard. The six that still need the helper -- the three stride refusals, the devComm clear, the context-destroy cleanup, and the unknown backend type -- keep it, with the justification moved next to them. setupWithBackend/setupAndFree move from GinHostTest into the suite that uses them, joined by a setup() on the public entry point. Verified: rccl-UnitTestsMicroGinHost 80 tests green in 154 ms; gin_host.cc coverage unchanged (100% lines, 74.43% branches); each moved test confirmed red by mutating its behaviour in the hipified tree. Co-authored-by: Pi Agent
…our they assert Every suite in gin-host-test.cc now follows the MICROTEST_README naming rule (GinHost<EntryPoint>Microtest) and every case follows Unit_Condition_ExpectedBehaviour. The pre-existing GinHostTest suite became the GinHostFixture base class, with its cases redistributed into GinHostProxyNthreadsMicrotest, GinHostProgressMicrotest and GinHostProgressThreadsMicrotest; the finalize-joins-the-workers case moved to GinHostFinalizeMicrotest, where it belonged. Cases whose name covered only part of what they asserted were split so that each name is honest about its own assertions: the two type-query accessors, their null-argument refusals, the signal-request success control, the env-type backend selection, and the traffic-class fallback. The file header no longer lists the suites: that list duplicated the code below it and MICROTEST_README.md, so each suite now carries its own banner and there is one place to keep current. Co-authored-by: Pi Agent
…-scope singleton FakeGin recovered itself from a file-scope singleton and ignored the ginInstance / listenComm / collComm / ginCtx handles production threads back into the plugin, so a handle mix-up passed silently. Each handle the fake hands out is now a tagged FakeHandle that names the fake and the role it was issued for; the vtable statics are thin trampolines that recover the fake from the handle and forward to a member function holding the behaviour. devices() and getProperties() take no handle in the plugin ABI, so those two alone still use the (renamed) handleless instance. A new GinHostFakeHandleMicrotest pins the detection: delivering a collComm to closeListen fails the test. Co-authored-by: Pi Agent
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
…omrccl-2642-gin-host-tests
… argument in the host fakes develop gave ncclSocketToString an explicit buffer-length parameter, but the rasnet and ras-peers test fakes still defined the old three-argument form, so the production callers in rasnet.cc and peers.cc linked against nothing. Match the new signature in both fakes and bound the snprintf with the caller's bufLen instead of the hardcoded SOCKET_NAME_MAXLEN + 1. Co-authored-by: Pi Agent
mch
marked this pull request as ready for review
October 14, 2026 16:20
…omrccl-2642-gin-host-tests
pvallem
approved these changes
Oct 14, 2026
Member
|
LGTM |
aduforest
approved these changes
Oct 14, 2026
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
ARGUS
Findings
5 more finding(s), not posted inline
projects/rccl/test/host/gin-host-test.cc:1487(Low) SetupFails_ClearsTheGinFields: the three stride assertions can never fail, since those fields are never poisonedprojects/rccl/test/host/gin-host-test.cc:1617(Low) DestroyContextFails test pins a pre-existing production leak of dc (gin_host.cc:486 returns before free at :488)projects/rccl/test/host/gin-host-test.cc:66(Low) Comment says ncclGinConnectOnce has only one AllGather; gin_host.cc:198 adds one per connectionprojects/rccl/test/host/group-test.cc:811(Low) SpinUntilAbortFlagObserved budget silently dropped from 5 s to waitUntil's 2 s defaultprojects/rccl/test/host/fakes/nccl_device_core_fakes.cc:21(Low) Comment says the ncclTeamLsa move is 'tracked separately'; README says it stays on purpose
This PR changes 2,359 lines across 23 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: LGTM once the ginCommCount = 7 poison at gin-host-test.cc:1032 is kept within NCCL_GIN_MAX_CONNECTIONS, so the fail: loop stops indexing past ginComms[4].
Comment @argus review to re-run, or @argus re-review after pushing.
…omrccl-2642-gin-host-tests
…proxy-test Merging develop brought in #12756, which defines a local ncclSetThreadName stub in rma-proxy-test.cc. This branch had already moved that symbol onto its owning fakes file (fakes/nccl_fakes.cc), which rccl-UnitTestsMicroRmaProxy also links, so the two definitions collided at link time: ld.lld: error: duplicate symbol: ncclSetThreadName(std::thread&, char const*, ...) Remove the local stub and note in the surrounding comment that symbols with an owning fakes file must not be redefined in this TU. Co-authored-by: Pi Agent
…stubs.cc gin-host-test.cc defined ncclTopoGetLocalGinDevs itself, under a comment claiming the symbol had no owning fakes file and that no topo fakes file could link into rccl-UnitTestsMicroGinHost. Neither is true: fakes/topo_stubs.cc is src/graph/topo.cc's fakes file, it is a floor of abort-on-call stubs and std::function seams with no heavy dependencies, and it links into this binary unchanged. Add the symbol to topo_stubs.cc as a controllable seam -- a seam rather than a plain abort because gin-host-test.cc drives both the device count, which sizes ginCommCount, and the failure arm. The default stays fail-loud, matching the rest of that file: this is the only binary that compiles a caller, and it installs its own behaviour from ResetGinHostGlobals, the same way it already installs bootstrap_stubs.cc's g_bootstrapAllGather. The hook now carries ASSERT_HOOK_MATCHES_PROD, which the hand-written definition could not have. No behaviour change: gin_host.cc coverage is identical (100% lines, 100% functions, 438/112 branches) and the full host suite passes. Co-authored-by: Pi Agent
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,359 lines across 23 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: Progress_PluginProgressFails_StoresTheStatusInAsyncResultAndExits (gin-host-test.cc:725) still passes with the error return NULL at gin_host.cc:82 deleted.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,381 lines across 25 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 ginCommCount = 7 poison at gin-host-test.cc:1039 still reads past ginComms[4], and the progress-fails, GIN_NCONNECTIONS and rail-handle tests still pass with the logic they target reverted.
Comment @argus review to re-run, or @argus re-review after pushing.
…read, and correct the coverage claim The affinity seams os_fakes.cc owns are reached from threads gin_host.cc spawns, so their mask vectors are written under a mutex. Until now the only thing stopping a reader racing them was a sentence in the header telling it to hold the lock. Add ncclOsCpuCountMasksSnapshot / ncclOsSetAffinityMasksSnapshot, which copy under the lock, and use them from gin-host-test.cc -- the one suite that spawns production threads. The single-threaded suites keep reading the vectors directly, which is safe for them and avoids churning init-test.cc. Joining a worker and reading its writes safely are separate problems; the comments now say so, because the previous wording let them blur together. Also replace the MICROTEST_README claim that every source-level branch in gin_host.cc is covered. Measured, it is 100% of lines and functions but 74% of branches, and the residue is not all NCCLCHECK/WARN as claimed: four INFO gates, two `while (1)` headers llvm-cov records an untakeable exit edge for, and 21 NCCLCHECK sites of which 13 have only the WARN gate left. The replacement states the groups and tells the reader how to re-derive them, including that lcov and llvm-cov do not agree on branch totals. Full host suite green; gin_host.cc coverage unchanged. Co-authored-by: Pi Agent
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,429 lines across 25 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 ginCommCount = 7 poison at gin-host-test.cc:1040 still reads past ginComms[4], and several tests still pass with the gin_host.cc logic they target mutated.
Comment @argus review to re-run, or @argus re-review after pushing.
7 exceeded NCCL_GIN_MAX_CONNECTIONS, so the fail path read past ginComms[4]. Co-authored-by: Pi Agent
Co-authored-by: Pi Agent
…eturns The test only waited on asyncResult, which a thread that stored the status and kept looping would also satisfy. Wait on the worker's own return and pin the progress call count at 1. Co-authored-by: Pi Agent
Every deregister case used one backend, where the slot index equals the connection index. Co-authored-by: Pi Agent
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,463 lines across 25 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: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
…omrccl-2642-gin-host-tests
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,517 lines across 25 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: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
…omrccl-2642-gin-host-tests Co-authored-by: Pi Agent
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,506 lines across 26 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: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
Contributor
pr-rocm-systems-13225-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-13225-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-13225-ruby64-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-13225-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 12 tests skipped. |
Contributor
pr-rocm-systems-13225-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
…omrccl-2642-gin-host-tests
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,495 lines across 26 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: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
1 task done
…omrccl-2642-gin-host-tests
dhanavikram
added a commit
that referenced
this pull request
Oct 28, 2026
…b helpers - gdr: the Hyper-V blocklist is one table that adds relaxed ordering on, the public API is checked to stop at the blocklist before its runtime probe, and the DMA-BUF probe table covers PD allocation and release failures (ScopedDmaBufProbe gains allocFails/deallocFails). The init-once PeerMem pair, covered by the public pair, is gone. - common: BaseCommInit starts from stale state, both BY_ID triggers and every prepost setting are driven, a failing resiliency init keeps the flush fd closed, and the speed and async-event tests cover a remote speed of 0, a local-only counter change and QP_REQ_ERR/QP_ACCESS_ERR. - init: an InfiniBand port defaults AR on, a successful relaxed-ordering probe enables it and IB_OOO_RQ keeps the AR threshold; ncclIbMakeVDevice is checked to wait for ncclIbMutex. Discovery runs on IbVerbsWorld.h, which FakeIbWorld now builds on. - gid-inspect covers the lower nqps clamp and both lookup filters; reg pins the page-boundary reuse check and the move of the last MR into a freed slot. - NetIbResetGlobals() (NetIbGlobals.h) runs in every net_ib TearDown; NetIbSupportSeams.h holds the gdr.cc probe trampolines; SysfsAnswers moves to os_fakes; nccl_fakes.cc takes #13225's weak ncclSetThreadName(std::thread&). - The single net_ib source-line pins become rows in _rccl_cast_pins.
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,493 lines across 25 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: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
…llNet teardown ResetDevRuntimeMicroFakes installs the HIP VMM emulator, which replaces g_hipMemGetAddressRange with a success stub. Running it after ResetHipFakes left that stub installed for the next suite in the binary, so P2pMicrotest.IpcRegisterBuffer_FreshRegistrationFailureClearsOutputs saw a working hipMemGetAddressRange and failed whenever --gtest_shuffle ordered a CollNet suite ahead of it. Co-authored-by: Pi Agent
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,495 lines across 26 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: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
…omrccl-2642-gin-host-tests
Contributor
There was a problem hiding this comment.
ARGUS
This PR changes 2,495 lines across 26 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: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
dhanavikram
added a commit
that referenced
this pull request
Oct 31, 2026
## Motivation
`src/transport/net_ib/{init,common,reg,gdr,net_ib_gid_inspect}.cc` and
the host helpers in `net_ib/common.h` implement IB discovery,
communicator setup, registration caching, GDR capability selection,
async-event handling, and fault-injection inspection. These host control
paths had no host-only coverage, leaving validation, cleanup, and
capability branches dependent on hardware-backed tests.
## Technical Details
Extends the `rccl-UnitTestsMicroTransport` binary introduced by PR
#13858. Tests mirror `src/transport/net_ib/` under
`test/host/transport/net_ib/` and cover:
- device discovery, PCI/NUMA normalization, virtual-device construction,
Data Direct, properties, and init/finalize lifecycle
- communicator defaults, receive posting, QP selection, load-balancing
weights, speed changes, GID refresh, and async events
- verbs and mlx5 MR registration, cache reuse/insertion/reference
release, DMA-BUF selection, and multi-device failure behavior
- peer-memory, runtime GDR, and DMA-BUF capability detection with
process-isolated once-cache checks
- device/communicator GID state, QP inspection, resiliency state,
recovery/probing GIDs, and QP error transitions
The binary reuses the existing transport infrastructure and `develop`'s
IB verbs, mlx5dv, GDR peermem/probe and socket fakes unchanged.
`os_fakes` gains `SysfsAnswers`, the bios_version/numa_balancing sysfs
script the gdr suite now shares with `net-ib-cast-gdr-test.cc`, and
`nccl_fakes.cc` carries #13225's weak no-op
`ncclSetThreadName(std::thread&, ...)`, which `init.cc` needs. It adds
the net_ib connect/p2p/resiliency/recovery owner seams and four shared
headers:
- `transport/net_ib/NetIbDmaBufProbe.h`, a scripted DMA-BUF probe shared
by the gdr and init suites;
- `transport/net_ib/NetIbGlobals.h`, whose `NetIbResetGlobals()` every
net_ib fixture's TearDown calls because the device tables are
process-wide;
- `transport/net_ib/NetIbSupportSeams.h`, the `gdr.cc` capability probes
a unit is renamed onto, since `gdr-test.cc` compiles the real `gdr.cc`;
- `IbVerbsWorld.h`, the fake HCAs and verbs libraries device discovery
runs against, which `FakeIbWorld` in `net-ib-cast-test-common.h` also
builds on.
`NetIb*` suites are added to the existing transport CTest category.
Runner, JUnit, and coverage-object registration are inherited from PR
#13858 because this deliberately uses the same binary; like the other
host-only microtests it stays out of `ci-precheckin.json` after #14381.
Two disabled forward specs document incidental production defects
without changing production code. The findings are filed under the
AICOMRCCL-2397 incidental-findings epic as AICOMRCCL-3112.
## Issue Tracking
JIRA ID: AICOMRCCL-2808
Incidental findings epic: AICOMRCCL-2397
## Stack
This PR is stacked on #13858 and targets
`users/dhanavikram/rma-socket-host-tests`.
## Test Plan
- Built and ran `rccl-UnitTestsMicroTransport` in standalone and in-tree
modes.
- Ran the complete host-only binary/guard pipeline.
- Ran shuffled/repeated transport tests under ASan, UBSan, LSan, and
TSan.
- Mutated five predicates/state updates in the hipified net_ib files and
confirmed the focused tests fail.
- Measured scoped `llvm-cov` coverage against all five hipified net_ib
`.cc` files and `common.h`.
- Ran four independent review angles: runtime safety, test
effectiveness, fake ownership, and stacked integration.
## Test Result
- All host-only binaries and guards pass; Transport: 237 enabled tests.
- Standalone and in-tree Transport targets pass.
- ASan/UBSan/LSan: 5 shuffled repetitions pass with leak detection
enabled.
- TSan: the complete Transport binary passes with no race/deadlock
reports in enabled test paths.
- 5 of 5 focused mutations caught.
- Review follow-up: dropping either errno check (`EOPNOTSUPP`,
`EPROTONOSUPPORT`) from `ncclMlx5dvDmaBufCapable` is caught.
- Second review round: 32 of 32 mutations caught, one or more per new
assertion.
- All 67 net_ib production functions are exercised; combined `.cc` line
coverage is 90.03%.
- Production findings are isolated under AICOMRCCL-2397; this PR changes
test files only.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
Made with [Cursor](https://cursor.com)
---------
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Dhanavikram Sekar <dhanavikram@users.noreply.github.com>
1 task done
Contributor
TheRock Submodule Bump ActivityNewest first
|
dlamd1dai
added a commit
that referenced
this pull request
Nov 4, 2026
…collectives (#14225) ## Motivation Split from AICOMRCCL-1999 (per-DevComm GIN backend selection, NCCL 2.31.2). Stacked on #14219 (AICOMRCCL-3418). Review and merge that PR first. This PR's base is its branch, so the diff shows only this change. The built-in GIN-SDMA alltoall and allreduce run only when Anvil SDMA is the communicator's default GIN backend: - `ncclAllToAllGinSdmaEligible` and the allreduce's `ginAllReduceBaseEligible` require `ncclGetGinType(comm) == ANVIL_SDMA`, and `ncclGetGinType` reports only `backends[0]`. - `ncclGinA2AInitOnce` and `ncclGinAllReduceInitOnce` leave `reqs.ginType` at `NCCL_GIN_TYPE_NONE`, so `ncclGinDevCommSetup` binds the first backend in the list, whatever it is. Once another backend can rank ahead of SDMA, these paths stop being chosen even though SDMA is active. Their kernels are compiled for SDMA only, so the DevComm has to bind SDMA by type rather than take the default. With the backend priorities in #14219, SDMA still ranks first whenever it is active, so no existing configuration changes behavior. This removes the assumption before anything depends on it. ## Technical Details - **`ncclGinTypeActive(comm, type)`** (`src/gin/gin_host.cc`, declared in `gin_host.h`): true if `type` is among `backends[0..numActiveBackends)`. It uses the same gates as `ncclGetGinType`: full GIN connectivity, and GIN state not finalized. - **Eligibility:** both SDMA collectives call `ncclGinTypeActive(comm, NCCL_GIN_TYPE_ANVIL_SDMA)` in place of the default-backend check. - **DevComm:** both `InitOnce` functions set `reqs.ginType = NCCL_GIN_TYPE_ANVIL_SDMA`. These internal DevComms are created with `NCCL_VERSION_CODE`, whose compat entry has no requirements filter, so the field reaches `ncclGinDevCommSetup` unchanged. Anvil SDMA reports strong and VA signal support, so an explicit request passes `ncclGinValidateSignalRequest` just as the default did. - **`ncclDevrWindowSorted`** moves from `dev_runtime.cc` into `dev_runtime.h`. This lets the alltoall eligibility test register a window and reach an eligible baseline. `test/host/wrap-test.cc` drops the private copy of this struct it had kept to match the layout. With SDMA as the only backend, or as the default, both paths behave exactly as before. ## Issue Tracking JIRA ID: AICOMRCCL-3419 ## Test Plan - `rccl-UnitTestsFixturesDebug` `GinAlltoAllEligibilityTest`: the mock comm now registers a window, so the baseline is eligible and each negative case fails if its gate is deleted. Before this change every case was rejected by the window lookup, so a deleted gate went unnoticed. New cases: - `SdmaActiveButNotDefault`: eligible with backends `{PROXY, SDMA}` and `{GDA, PROXY, SDMA}`. - `SdmaNotActive`: not eligible with `{GDA}`, `{PROXY}` or `{GDA, PROXY}`. - `GinFinalized`: not eligible. `MockBackendSlotMatchesGinTypeAccessor` is removed, since the baseline now covers it. - `rccl-UnitTestsMicroGinHost` (`test/host/gin-host-test.cc`, on the fixtures from #13225): - `GinHostTypeQueryMicrotest.GinTypeActive_TypeInANonDefaultSlot_ReportsActive` and `GinTypeActive_RailOnlyFinalizedOrNullComm_ReportsInactive` cover the helper. - `GinHostBackendSelectMicrotest.DevCommSetup_RequestedTypeIsNotTheDefault_BindsThatBackend`: with `{PROXY, SDMA}`, a DevComm requesting SDMA records `backendIndex == 1`, and one without a type records 0. - `rccl-UnitTestsMicro`, which compiles `wrap-test.cc` and `dev_runtime.cc` against the moved struct. - The allreduce eligibility check cannot reach its backend gate in a host test: it first calls `isSymmetricKernelRequested`, which runs `ncclSymkInitOnce`. It uses the same helper as alltoall. - Mutation check: undo each part of the change and confirm the matching new test fails. ## Test Result `ctr2-alola-compile-12` (build) and MI300X `ctr-cx63-mi300x-12` (run). Debug build with `ENABLE_ROCSHMEM_GIN=ON`, `GPU_TARGETS=gfx942`, amdclang++ from ROCm nightly 10.3.0a20261024: ``` rccl-UnitTestsFixturesDebug GinAlltoAll/AllReduceEligibilityTest [ PASSED ] 21 tests (2 skipped: param-pinned cases) same, NCCL_GIN_A2A_ENABLE=0 [ PASSED ] 12 tests (A2ADisabledByParam runs) same, NCCL_GIN_A2A_MIN_BYTES=64 MiB [ PASSED ] 12 tests (BelowMinBytesParam runs) rccl-UnitTestsMicroGinHost [ PASSED ] 20 tests rccl-UnitTestsMicro (full) [ PASSED ] 2213 tests ``` After the rebase onto #14219 (`137f0eb79f0`, which brings in #13225's rewrite of `gin-host-test.cc`), host tests in the CI image `rocm/dev-ubuntu-22.04:7.2`, shuffled. The GPU results above are from the earlier tree: ``` rccl-UnitTestsMicroGinHost [ PASSED ] 96 tests rccl-UnitTestsMicro (full) [ PASSED ] 3002 tests ``` Each mutation below made the matching new test fail. With the change restored, all pass: | Mutation | Failing test | |---|---| | Old alltoall gate (`ncclGetGinType == ANVIL_SDMA`) | `GinAlltoAllEligibilityTest.SdmaActiveButNotDefault` (both backend orders) | | `ncclGinTypeActive` reads only `backends[0]` | `GinHostTypeQueryMicrotest.GinTypeActive_TypeInANonDefaultSlot_ReportsActive` (re-run after the rebase) | No unit test covers the `reqs.ginType` lines in the two `InitOnce` functions, which create a real DevComm. `DevCommSetup_RequestedTypeIsNotTheDefault_BindsThatBackend` shows that `ncclGinDevCommSetup` honors that field. A multi-GPU run cannot exercise the non-default case yet, because SDMA always ranks first when it is active. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
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
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.
Brings
src/gin/gin_host.ccto full host-only test coverage inrccl-UnitTestsMicroGinHost, and tidies the surrounding microtest infrastructure the work exposed.PR #12475 added the first host-only coverage of this file (the GIN proxy NUMA affinity pin) and PR #12245 the
GIN_PROXY_NTHREADSprogress-thread behaviour, which together left it at 59% of lines and 34% of branches. This builds out the rest: every function in the file now has a test, across nine new suites (one per entry point) covering type query, signal requests, connect-once, backend select, devComm setup/free, finalize, register and last-error query. Coverage ofgin_host.ccgoes from 59.01% to 100% of lines, 66.67% to 100% of functions, and 33.79% to 74.43% of branches; no source-level branch is left uncovered, the remainder beingNCCLCHECK/WARNmacro expansions at call sites whose failure arm no test drives.The test TU's
ncclGin_tfake gains a uniformScriptedFailureper plugin entry point ("fail call N of this function") plus argument recorders, so each failure arm is reachable without a bespoke fake per test.FakeGinhands out tagged handles rather than relying on a file-scope singleton, so tests that hold more than one context stay independent.Four reusable utilities that had accumulated inside the test TU are promoted to shared headers for the other host-only suites:
ScopeExit.h,ScriptedFailure.h,WaitUntil.handJoinThreadsOnExit.h.Externals the suite had been defining locally move onto the fakes file named for the production TU that defines each, per
MICROTEST_README.md:bootstrapAllGathertofakes/bootstrap_stubs.cc,ncclOsCpuCount/ncclOsSetAffinitytofakes/os_fakes.cc,ncclSetThreadNametofakes/nccl_fakes.cc, and thencclTeam*accessors to a newfakes/nccl_device_core_fakes.cc. Each moved seam picks upASSERT_HOOK_MATCHES_PRODsignature-drift checking from its new home.Testing:
run_host_tests.sh allpasses — all 13 host-only binaries green, 87GinHosttests, shuffled. Ten mutations of the hipifiedgin_host.cc(dropping the NULL-window check, the multi-segment DMABUF requirement, the stride-multiple check, the empty-slot skip, the early return on a reported error, the context round-up, the connectivity check inncclGetGinType, the VA-signal requirement check, the devComm reset on failure, and the fail-pathcloseListen) each turn exactly the matching test red.JIRA ID : AICOMRCCL-2642