Skip to content
Merged
Show file tree
Hide file tree
Changes from 14 commits
Commits
Show all changes
31 commits
Select commit Hold shift + click to select a range
e235402
[AICOMRCCL-2642] [rccl] Cover every gin_host.cc entry point with host…
mch Oct 9, 2026
5a728ca
[AICOMRCCL-2642] [rccl] Move gin-host-test fakes onto their owning fa…
mch Oct 9, 2026
2bee9cb
[AICOMRCCL-2642] [rccl] Promote reusable gin-host-test utilities to s…
mch Oct 9, 2026
6ecbf53
[AICOMRCCL-2642] [rccl] Pin the contracts gin-host-test's thin tests …
mch Oct 9, 2026
236091a
[AICOMRCCL-2642] [rccl] Drive the gin-host devComm setup tests throug…
mch Oct 9, 2026
261597c
[AICOMRCCL-2642] [rccl] Name the gin-host microtests after the behavi…
mch Oct 9, 2026
7e2749b
[AICOMRCCL-2642] [rccl] Give FakeGin tagged handles instead of a file…
mch Oct 9, 2026
1424d45
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 13, 2026
e71d114
[AICOMRCCL-2642] [rccl] Follow ncclSocketToString's new buffer-length…
mch Oct 13, 2026
807d2a3
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 14, 2026
f249eb4
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 16, 2026
4ceee3a
[AICOMRCCL-2642] [rccl] Drop duplicate ncclSetThreadName stub in rma-…
mch Oct 16, 2026
70b8e91
[AICOMRCCL-2642] [rccl] Move ncclTopoGetLocalGinDevs onto fakes/topo_…
mch Oct 16, 2026
3b3e25e
[AICOMRCCL-2642] [rccl] Make the os_fakes affinity recorders safe to …
mch Oct 16, 2026
0f8108e
[AICOMRCCL-2642] [rccl] Use an in-range ginCommCount poison
mch Oct 16, 2026
55d89f5
[AICOMRCCL-2642] [rccl] Trim the host-test helper comments
mch Oct 16, 2026
7723c90
[AICOMRCCL-2642] [rccl] Assert the failing progress thread actually r…
mch Oct 16, 2026
451bb55
[AICOMRCCL-2642] [rccl] Cover GIN_NCONNECTIONS on its own
mch Oct 16, 2026
1a9ad14
[AICOMRCCL-2642] [rccl] Assert the handle array the rail-only team co…
mch Oct 16, 2026
06e1557
[AICOMRCCL-2642] [rccl] Assert which device each GIN connection is op…
mch Oct 16, 2026
c73c279
[AICOMRCCL-2642] [rccl] Assert the stride reciprocal and device types…
mch Oct 16, 2026
62c5e77
[AICOMRCCL-2642] [rccl] Cover deregistration across two active backends
mch Oct 16, 2026
198b0af
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 22, 2026
8e56909
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 26, 2026
a46fffe
Merge branch 'develop' into users/mahughes/aicomrccl-2642-gin-host-tests
mch Oct 26, 2026
de98a72
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 27, 2026
c66f9f8
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 28, 2026
f7024a7
[AICOMRCCL-2642] [rccl] Reset HIP fakes after dev-runtime fakes in Co…
mch Oct 28, 2026
ecb9c36
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 29, 2026
51df229
Merge remote-tracking branch 'origin/develop' into users/mahughes/aic…
mch Oct 29, 2026
60291ba
Merge branch 'develop' into users/mahughes/aicomrccl-2642-gin-host-tests
mch Oct 30, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions projects/rccl/test/host/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -116,6 +116,7 @@ function(rccl_define_micro_source_lists)
# real IsArchMatch, for wrap-test.cc (path is build-directory-relative).
${RCCL_HOST_HIPIFY_DIR}/src/misc/archinfo.cc
fakes/dev_runtime_micro_fakes.cc
fakes/nccl_device_core_fakes.cc # src/nccl_device/core.cc ncclTeam* accessors
fakes/ibvwrap_fakes.cc # src/misc/ibvwrap.cc PD / MR registration seams
fakes/nccl_fakes.cc
fakes/libc_fakes.cc
Expand Down Expand Up @@ -395,6 +396,10 @@ function(rccl_define_micro_source_lists)
fakes/nccl_fakes.cc
fakes/hip_fakes.cc
fakes/env_fakes.cc
fakes/bootstrap_stubs.cc # src/bootstrap.cc AllGather seam
fakes/nccl_device_core_fakes.cc # src/nccl_device/core.cc ncclTeam* accessors
fakes/os_fakes.cc # src/os/linux.cc proxy-thread affinity seams
fakes/topo_stubs.cc # src/graph/topo.cc ncclTopoGetLocalGinDevs seam
../common/main_altrsmi.cpp
PARENT_SCOPE)

Expand All @@ -416,6 +421,7 @@ function(rccl_define_micro_source_lists)
fakes/diagnostics_p2p_device_fakes.cc # src/diagnostics/device/p2p.cu
fakes/hip_fakes.cc
fakes/libc_fakes.cc
fakes/nccl_device_core_fakes.cc # src/nccl_device/core.cc ncclTeam* accessors
fakes/nccl_fakes.cc
fakes/topo_stubs.cc
fakes/transport_net_ib_fakes.cc # src/transport/net_ib/connect.cc params
Expand Down
68 changes: 68 additions & 0 deletions projects/rccl/test/host/JoinThreadsOnExit.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,68 @@
/*************************************************************************
* Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/

#ifndef RCCL_TEST_HOST_JOIN_THREADS_ON_EXIT_H_
#define RCCL_TEST_HOST_JOIN_THREADS_ON_EXIT_H_

#include <functional>
#include <thread>
#include <utility>
#include <vector>

// JoinThreadsOnExit -- spawn worker threads, stop and join them on scope exit
//
// Destroying a still-joinable std::thread calls std::terminate(), so a test that
// spawns a worker and then trips a fatal `ASSERT_*` would abort the whole binary
// rather than fail one test. Spawning through this type means the stop-and-join
// happens on every exit path, including that one.
//
// The loop a worker runs is usually ended by a flag the unit under test owns
// rather than by the thread returning on its own, so the stop action is a
// callback supplied by the caller -- this type knows nothing about what makes
// the workers finish, only that it must be done before joining them. Pass no
// callback for workers that finish by themselves.
//
// Usage:
//
// JoinThreadsOnExit workers([gs] { gs->stopSignal.store(true); });
// workers.spawn([gs] { ncclGinProgress(gs, 0); });
// ASSERT_TRUE(waitUntil([&] { return fake_.progressCalls.load() > 0; }));
// // workers' destructor raises the stop signal and joins, even on the ASSERT.
//
// Non-copyable and non-movable so the threads have exactly one owner; C++17
// guaranteed copy elision still allows a factory that binds the stop callback to
// return one by value (`auto workers = progressWorkers(gs);`).
class JoinThreadsOnExit {
public:
explicit JoinThreadsOnExit(std::function<void()> stop = {}) : stop_(std::move(stop)) {}
~JoinThreadsOnExit() { stopAndJoin(); }

JoinThreadsOnExit(const JoinThreadsOnExit&) = delete;
JoinThreadsOnExit& operator=(const JoinThreadsOnExit&) = delete;
JoinThreadsOnExit(JoinThreadsOnExit&&) = delete;
JoinThreadsOnExit& operator=(JoinThreadsOnExit&&) = delete;

template <class Fn>
void spawn(Fn&& fn) { threads_.emplace_back(std::forward<Fn>(fn)); }

// Also callable before the scope ends, for a test that asserts on what the
// workers left behind once they have all stopped. Idempotent.
void stopAndJoin()
{
if (threads_.empty()) return;
if (stop_) stop_();
for (auto& t : threads_) {
if (t.joinable()) t.join();
}
threads_.clear();
}

private:
std::function<void()> stop_;
std::vector<std::thread> threads_;
};

#endif // RCCL_TEST_HOST_JOIN_THREADS_ON_EXIT_H_
74 changes: 70 additions & 4 deletions projects/rccl/test/host/MICROTEST_README.md
Original file line number Diff line number Diff line change
Expand Up @@ -318,10 +318,33 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary:
`ENABLE_WARP_SPEED` onto this binary; the scenes leave `warpSpeedEnabled`
false, so the stubs exist to link. See `test_categories_micro_connect.yaml`.
- **`rccl-UnitTestsMicroGinHost`** — `gin/gin_host.cc` (`GIN_HOST_CC_PATH`, from
`gin-host-test.cc`); suite `GinHostTest.*`. NVIDIA/nccl#2279
`NCCL_GIN_PROXY_NTHREADS` progress-thread assignment. Its own binary, not
`gin-host-test.cc`); suites `GinHost<EntryPoint>Microtest`, one per entry
point. Covers every function in the file: type
negotiation, signal-requirement validation, `ncclGinConnectOnce` (including
NVIDIA/nccl#2279 `NCCL_GIN_PROXY_NTHREADS` progress-thread assignment and the
strided team a rail-only comm connects), devComm setup (per-backend version
tables, stride validation, failure cleanup) and free, finalize, window
registration, and last-error queries, plus the proxy-thread CPU affinity pin.
A scripted `ncclGin_t` in the test file stands in for the plugin, so every
failure arm is reachable by failing a chosen call. Its own binary, not
sharing `rccl-UnitTestsMicro`: `gin-plugin-init-test.cc` already defines
`ncclParamGinEnable` there. See `test_categories_micro_gin_host.yaml`.
Every line and every function in `gin_host.cc` is covered (100% of both);
branch coverage is 74%. The residue is in three groups, none of them a
source-level decision the tests skip:
- the debug-level gate inside `INFO` / `INFO_LOC`, which host tests run with
logging off (4 call sites);
- the `while (1)` headers of the two progress loops, where `llvm-cov` records
an exit edge a literal `1` can never take (2 sites);
- `NCCLCHECK`-family call sites (21). At 13 of them the check's own failure
arm *is* driven and only the `WARN` inside the macro is left; at the other
8 the failure arm is undriven, and five of those are `ncclCalloc` sites
that would need an allocation injector this binary does not have.

Re-derive rather than trusting those numbers: `BRDA` records for the file in
`build/coverage/rccl-UnitTestsMicroGinHost/coverage.lcov` give the per-line
breakdown, and note that lcov and `llvm-cov` count branches differently, so
the two totals will not agree.

- **`rccl-UnitTestsMicroRmaProxy`** — `src/rma/rma_proxy.cc` (via
`RMA_PROXY_CC_PATH`, from `rma-proxy-test.cc`); suite
Expand Down Expand Up @@ -424,7 +447,9 @@ test:
`ScopedHook` helper in `ScopedHook.h` (see
[Installing per-test behaviour with `ScopedHook`](#installing-per-test-behaviour-with-scopedhook)),
which installs the hook, counts calls, and restores the previous behaviour
automatically on scope exit.
automatically on scope exit. Reach for the other
[shared test-side helpers](#shared-test-side-helpers) before hand-rolling a
failure counter, a wait loop, or a scope guard.
4. **Only exercise faked seams.** Every external symbol the `#include`d `.cc`
reaches must be satisfied by `fakes/`: a missing symbol surfaces as a
link error, a wrongly
Expand Down Expand Up @@ -476,6 +501,7 @@ symbol.
| `src/group.cc` | `fakes/group_fakes.cc` |
| `src/config/algorithm_*.cc` | compiled real (no fakes file) |
| `src/config/collconfig.cc` (targets that do not compile the real file) | `fakes/collconfig_fakes.cc` |
| `src/debug.cc` (`ncclDebugLog`, `ncclDebugLevelMask`, `ncclSetThreadName`, ...) | `fakes/nccl_fakes.cc` |
| `src/enqueue/enqueue.cc`'s own symbols (targets that do not compile the real file) | `fakes/enqueue_fakes.cc` |
| `src/init.cc` comm lifecycle + CTA/channel params | `fakes/comm_fakes.cc` |
| `src/init_nvtx.cc` | `fakes/init_nvtx_fakes.cc` |
Expand All @@ -489,6 +515,7 @@ symbol.
| `src/misc/rocmwrap.cc` | `fakes/rocmwrap_fakes.cc` |
| `src/misc/strongstream.cc` | `fakes/strongstream_stubs.cc` |
| `src/misc/utils.cc` | `fakes/utils_fakes.cc` |
| `src/nccl_device/core.cc` host `ncclTeam*` accessors | `fakes/nccl_device_core_fakes.cc` (except `ncclTeamLsa`; see below) |
| `src/os/linux.cc` | `fakes/os_fakes.cc` |
| `src/plugin/env.cc` | `fakes/env_plugin_fakes.cc` |
| `src/plugin/gin.cc`, `src/gin/gin_host.cc` (targets that do not compile the real file) | `fakes/gin_fakes.cc` |
Expand Down Expand Up @@ -519,13 +546,24 @@ owns one driven seam, and `os_fakes.cc` holds working implementations with no se
read the suffix as a guarantee; read the file's header comment, which states what it owns. If you
add a file, pick the suffix matching its majority content and say so at the top.

**Recorders a production thread can reach.** Most fakes are touched only by the test's own thread,
so their counters and recording vectors are plain `int`s and `std::vector`s. A fake reached from a
thread the unit under test spawns needs more: `os_fakes.cc`'s affinity pair is the current example
(`gin_host.cc`'s proxy progress workers call `ncclOsCpuCount` / `ncclOsSetAffinity` from their own
threads), so its counter is a `std::atomic<int>` and its mask vectors are written under
`g_ncclOsAffinityMutex` and read through `ncclOs*MasksSnapshot()`. If you seam something a
spawned thread reaches, do the same, and keep a direct-read path for the single-threaded suites
rather than churning them. Note that joining the worker and reading safely are two different
problems: the join is what makes the write have happened, the snapshot is what makes the read
safe.

**`// UNDRIVEN`.** A seam carrying this marker is declared so the binary links and so an accidental
call is visible, NOT because its path is covered. It is a link-floor entry with a chosen default,
and that default silently selects which production arm runs. Driving a seam means deleting its
marker. The marker travels with the declaration rather than a block comment so it cannot drift from
what it describes. Call *counters* do not take the marker unless the counter itself is unread.

Five things do NOT follow the TU-per-file rule, deliberately:
These do NOT follow the TU-per-file rule, deliberately:

- `fakes/collective_stubs.cc` is a fail-loud floor for the collective *launch*
pipeline (`ncclLaunchKernel` and friends), which `enqueue.cc` itself defines.
Expand All @@ -551,6 +589,12 @@ Five things do NOT follow the TU-per-file rule, deliberately:
- `rcclParamIntraGraphGen` stays in `fakes/init_fakes.cc` because its owner
(`graph/rccl_graph_gen.cc:34`) has no fakes file at all. `rcclEffectiveP2pBatchEnable`
did have one and moved to `fakes/enqueue_fakes.cc`.
- `ncclTeamLsa` stays in `fakes/devcomm_fakes.cc` rather than moving to
`fakes/nccl_device_core_fakes.cc` with the other `src/nccl_device/core.cc`
accessors: its default silently satisfies the `nRanks ==
comm->nRanks` comparison every `devcomm-test.cc` filter makes, and the trap
comment warning about that is what keeps those tests honest. Moving the hook
without moving the readers that depend on the warning is how it would get lost.
- `IsArchMatch` and the `allocTracker` data symbol stay in `p2p-test.cc`
itself rather than a fakes file, because neither has an owning production
TU to name a fakes file after: `IsArchMatch` is declared in the
Expand Down Expand Up @@ -707,6 +751,28 @@ test can't know which at authoring time), pair the factories that drive each
version's arm so the test passes regardless of the toolchain — see
`ForceLegacyCudaRegister` + `ForceLegacyIpcCapable` in `p2p-test.cc`.

## Shared test-side helpers

The scaffolding a test needs on *its* side of the seam — scripting a fake's
failure, waiting for a worker thread, cleaning up on the fatal-assert path — is
the same in every binary, so it lives in a header next to the tests rather than
being re-invented per TU. Each is one small header named after the type it
declares:

| Header | Use it for |
|---|---|
| `ScopedHook.h` | Install behaviour on a `std::function` seam, count the calls, restore the previous behaviour on scope exit. See [Installing per-test behaviour with `ScopedHook`](#installing-per-test-behaviour-with-scopedhook). |
| `ScriptedFailure.h` | "Fail call N of this entry point." A fake holds one per entry point and asks `at(++calls)` what to return; the test writes `fake_.failConnect = {ncclSystemError, 2}`. Replaces the `g_fooCalls` + `g_failFooCall` pair that otherwise grows per seam. |
| `WaitUntil.h` | `waitUntil(pred)` — bounded poll for something another thread will do, returning false when the budget elapses so the test fails with its own message instead of hanging. Never `sleep_for` and hope. |
| `JoinThreadsOnExit.h` | Spawn worker threads and stop/join them on every exit from the scope. The stop action is a callback, because what ends the worker's loop belongs to the unit under test. |
| `ScopeExit.h` | The one-off cleanup that has no purpose-built RAII type. A fatal `ASSERT_*` returns immediately, so cleanup written after it does not run. |

When a test grows a helper that is not specific to its unit under test, promote
it here rather than leaving it in the TU — a per-TU copy is how four tests end
up with four slightly different "wait for the thread" loops. Helpers that *are*
specific to one unit stay with it (`TaskPrepScene.h` for the task-prep scene,
`fakes/ras_*_test_support.h` for the RAS fixtures).

## Dealing with each kind of dependency

When the link fails with `undefined symbol: foo`, find `foo` and
Expand Down
52 changes: 52 additions & 0 deletions projects/rccl/test/host/ScopeExit.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,52 @@
/*************************************************************************
* Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/

#ifndef RCCL_TEST_HOST_SCOPE_EXIT_H_
#define RCCL_TEST_HOST_SCOPE_EXIT_H_

#include <utility>

// ScopeExit -- run this cleanup on every exit from the scope
//
// A fatal assertion (`ASSERT_*`) returns from the test body immediately, so
// cleanup written after the assertion simply does not run -- which in a host
// test usually means a leaked allocation, a still-joinable std::thread (that
// calls std::terminate from its destructor), or a fake left pointing at a dead
// stack local, so one failing test takes the rest of the binary with it. Put
// that cleanup in a ScopeExit at the point the resource is acquired and it runs
// on the fatal path too.
//
// Usage (CTAD deduces the callable's type):
//
// std::thread worker([&] { ... });
// ScopeExit joinWorker([&] { stop.store(true); worker.join(); });
// ASSERT_TRUE(waitUntil([&] { return started.load(); })); // safe to bail here
//
// Prefer a purpose-built RAII type where one already exists (ScopedHook for a
// std::function seam, JoinThreadsOnExit for a set of worker threads); ScopeExit
// is for the one-off cleanup that does not have one.
//
// Non-copyable and non-movable so the cleanup cannot run twice; C++17 guaranteed
// copy elision still allows a factory to return one by value.
template <typename Fn>
class ScopeExit {
public:
explicit ScopeExit(Fn fn) : fn_(std::move(fn)) {}
~ScopeExit() { fn_(); }

ScopeExit(const ScopeExit&) = delete;
ScopeExit& operator=(const ScopeExit&) = delete;
ScopeExit(ScopeExit&&) = delete;
ScopeExit& operator=(ScopeExit&&) = delete;

private:
Fn fn_;
};

template <typename Fn>
ScopeExit(Fn) -> ScopeExit<Fn>;

#endif // RCCL_TEST_HOST_SCOPE_EXIT_H_
56 changes: 56 additions & 0 deletions projects/rccl/test/host/ScriptedFailure.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,56 @@
/*************************************************************************
* Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/

#ifndef RCCL_TEST_HOST_SCRIPTED_FAILURE_H_
#define RCCL_TEST_HOST_SCRIPTED_FAILURE_H_

#include "nccl.h"

// ScriptedFailure -- "fail call N of this entry point"
Comment thread
mch marked this conversation as resolved.
Resolved
//
// A test that needs the *second* connect() to fail otherwise grows a bespoke
// pair of file-scope knobs per seam (`g_fooCalls` + `g_failFooCall`) and a
// hand-written comparison inside every fake. ScriptedFailure is that pattern
// once: the fake asks what to return for this call number, and the test says
// which call fails with what.
//
// `onCall` is 1-based; 0 means "every call", so a test that only cares that the
// seam fails does not have to know how many times production calls it.
//
// Usage -- in the fake, next to the counter the test asserts on:
//
// int connectCalls = 0;
// ScriptedFailure failConnect; // default: every call succeeds
//
// ncclResult_t Connect(...) {
// ncclResult_t ret = failConnect.at(++connectCalls);
// if (ret != ncclSuccess) return ret;
// ...succeed...
// }
//
// and in the test:
//
// fake_.failConnect = {ncclSystemError, 2}; // the second connect fails
//
// Deliberately a plain aggregate with no counter of its own: tests assign whole
// scripts with `= {ncclSystemError, 2}` and reset them with `= {}`, and the call
// counter usually has to be visible to the test's own assertions anyway. Keeping
// the count outside also leaves the fake free to share one script across several
// entry points, or to count calls the script does not care about.
struct ScriptedFailure {
ncclResult_t result = ncclSuccess;
int onCall = 0;

// Status the seam should return for call number `callNumber` (1-based).
ncclResult_t at(int callNumber) const
{
if (result == ncclSuccess) return ncclSuccess;
if (onCall == 0 || onCall == callNumber) return result;
return ncclSuccess;
}
};

#endif // RCCL_TEST_HOST_SCRIPTED_FAILURE_H_
40 changes: 40 additions & 0 deletions projects/rccl/test/host/WaitUntil.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,40 @@
/*************************************************************************
* Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/

#ifndef RCCL_TEST_HOST_WAIT_UNTIL_H_
#define RCCL_TEST_HOST_WAIT_UNTIL_H_

#include <chrono>
#include <thread>

// waitUntil -- bounded poll for something another thread will do
//
// A test that spawns a worker cannot assert on the worker's effect immediately,
// and must not sleep-and-hope either: a fixed sleep is both slower than it needs
// to be on a good run and still flaky on a loaded machine. waitUntil returns as
// soon as `pred()` holds, and returns false once `budget` has elapsed so the
// test fails with its own message instead of hanging until the suite's timeout:
//
// ASSERT_TRUE(waitUntil([&] { return fake_.progressCalls.load() > 0; }))
// << "the progress thread never called ginProgress";
//
// The default budget is deliberately generous: it is never waited for on a
// passing run (the predicate ends the wait), and it only has to be long enough
// that a heavily loaded CI machine does not trip it.
inline constexpr auto kWaitUntilBudget = std::chrono::milliseconds(2000);

template <typename Pred>
bool waitUntil(Pred pred, std::chrono::milliseconds budget = kWaitUntilBudget)
{
const auto deadline = std::chrono::steady_clock::now() + budget;
while (!pred()) {
if (std::chrono::steady_clock::now() >= deadline) return false;
std::this_thread::yield();
}
return true;
}

#endif // RCCL_TEST_HOST_WAIT_UNTIL_H_
Loading
Loading