mch · Oct 31, 2026 · Oct 9, 2026 · Oct 9, 2026 · Oct 9, 2026 · Oct 9, 2026
@@ -239,6 +239,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
@@ -623,6 +624,7 @@ function(rccl_define_micro_source_lists)
     fakes/hip_fakes.cc
     fakes/env_fakes.cc
     fakes/devcomm_fakes.cc
+    fakes/nccl_device_core_fakes.cc      # ncclTeamRail, read by the GIN cost model
     fakes/os_fakes.cc
     fakes/profiler_fakes.cc
     fakes/strongstream_stubs.cc
@@ -803,10 +805,13 @@ function(rccl_define_micro_source_lists)
   # rccl-UnitTestsMicro, and gin_fakes.cc supplies ncclGinQueryLastError to Init.
   set(TEST_MICRO_GIN_HOST_SOURCE_FILES
     gin-host-test.cc
-    fakes/devcomm_fakes.cc               # ncclTeamRail
     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)
 
@@ -891,6 +896,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

@@ -0,0 +1,50 @@
+/*************************************************************************
+ * 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>
+
+// Spawn worker threads; stop and join them on every scope exit, including a fatal
+// ASSERT_* (destroying a joinable std::thread would std::terminate the binary).
+// The stop action is a caller-supplied callback because the worker's loop is ended
+// by a flag the unit under test owns; omit it for workers that finish by themselves.
+// Non-copyable and non-movable so the threads have exactly one owner.
+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_
@@ -417,10 +417,33 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary:
   because `misc/param.cc` defines `ncclGetEnv`, `initEnv` and `ncclLoadParam`, which
   the shared fakes in `rccl-UnitTestsMicro` also define.
 - **`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-UnitTestsMicroPlugin`** — `plugin/net.cc`, `plugin/plugin_open.cc` and the
   `plugin/net/net_v6.cc` to `net_v12.cc` shims (via `PLUGIN_NET_CC_PATH`,
   `PLUGIN_OPEN_CC_PATH` and `NET_V<N>_CC_PATH`, one test TU each); suite `Plugin*`. Its
@@ -636,7 +659,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.
    A death test takes `DEATH_BY_SEGV`, `DisableCoreDumpsForExpectedCrash()` and
    `ScopedDeathTestStyle` from `ExpectedCrash.h`. If the crash needs a NULL, hide it
    from the optimizer with `asm volatile("" : "+r"(p))`, or it may drop the crash.
@@ -697,6 +722,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, `allocTracker` | `fakes/comm_fakes.cc`; `rccl-UnitTestsMicroRmaProxy` keeps a local `allocTracker` because that isolated target does not link the broader comm fake |
 | `src/init_nvtx.cc` | `fakes/init_nvtx_fakes.cc` |
@@ -716,6 +742,7 @@ symbol.
 | `src/misc/socket.cc` | `fakes/socket_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/profiler.cc` (`ncclProfilerProxyTraceDumpIfAny`, recorded) | `fakes/profiler_fakes.cc` |
@@ -754,13 +781,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.
@@ -786,6 +824,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.
 
 An aggregation header includes the per-TU headers a unit's tests use and
 declares the reset that chains their per-TU resets, and defines no seams itself.
@@ -936,6 +980,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

@@ -0,0 +1,35 @@
+/*************************************************************************
+ * 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>
+
+// Run this cleanup on every exit from the scope, including a fatal ASSERT_* (which
+// returns before any cleanup written after it, leaking the resource or leaving a
+// fake pointing at a dead stack local). Prefer a purpose-built RAII type where one
+// exists (ScopedHook, JoinThreadsOnExit); ScopeExit is for the one-off.
+// Non-copyable and non-movable so the cleanup cannot run twice.
+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_
@@ -0,0 +1,29 @@
+/*************************************************************************
+ * 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"
+
+// "Fail call N of this seam": the fake asks at(++calls); onCall is 1-based, 0 fails
+// every call. A plain aggregate with no counter of its own, so tests can assign
+// (`= {ncclSystemError, 2}`) and reset (`= {}`) whole scripts, and one script can be
+// shared across seams whose counters the test asserts on anyway.
+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_
@@ -0,0 +1,30 @@
+/*************************************************************************
+ * 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>
+
+// Bounded poll for something another thread will do: returns as soon as pred() holds,
+// or false once budget elapses so the test fails with its own message rather than
+// hanging. The default budget is generous because it is never waited for on a passing
+// run, and only has to outlast a heavily loaded CI machine.
+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_
@@ -17,6 +17,7 @@
 #include "fakes/hip_fakes.h"      // the shared g_hip* seams this suite drives
 #include "fakes/nccl_fakes.h"     // g_ncclProxyClientGetFdBlocking
 #include "fakes/devcomm_fakes.h"  // g_ncclTeamLsa
+#include "fakes/nccl_device_core_fakes.h"  // g_ncclTeamWorld and the rest of the ncclTeam* family
 
 // param.h's NCCL_PARAM caches its value in a function-local static, so a param
 // read once is frozen for the process. fakes/param_redirect.h -- shared with
@@ -5379,7 +5380,7 @@ class DevrWorldToLsaRankSymTest : public DevrWorldToLsaRankTest {
     // descriptors carry *our own* position in that team -- the member test is
     // relative to it -- so world rank 4 is lsa rank 0, and a returned index is
     // plainly distinguishable from the world rank it came from.
-    g_devrTeamWorld = [](ncclComm_t) { return ncclTeam_t{8, 4, 1}; };
+    g_ncclTeamWorld = [](ncclComm_t) { return ncclTeam_t{8, 4, 1}; };
     g_ncclTeamLsa   = [](ncclComm_t) { return ncclTeam_t{4, 0, 1}; };
   }
   void TearDown() override {

@@ -14,6 +14,7 @@
 #include "bootstrap.h"
 
 #include "fakes/bootstrap_stubs.h"
+#include "fakes/signature-drift.h"
 
 bool g_bootstrapNetInitFail = false;
 ncclResult_t bootstrapNetInit() { return g_bootstrapNetInitFail ? ncclSystemError : ncclSuccess; }
@@ -25,6 +26,7 @@ std::function<ncclResult_t(void* commState, void* allData, int size)> g_bootstra
 ncclResult_t bootstrapAllGather(void* commState, void* allData, int size) {
   return g_bootstrapAllGather(commState, allData, size);
 }
+ASSERT_HOOK_MATCHES_PROD(g_bootstrapAllGather, bootstrapAllGather);
 // Only init.cc:2958 reaches this, to hand the rocSHMEM unique id to the other
 // ranks. Succeeds without touching bcastData; no test reaches it yet.
 ncclResult_t bootstrapBroadcast(void*, int, int, int, void*, int) { return ncclSuccess; }