Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
4 changes: 2 additions & 2 deletions projects/rccl/test/host/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -289,7 +289,7 @@ function(rccl_define_micro_source_lists)
../common/main_altrsmi.cpp
PARENT_SCOPE)

# rccl-UnitTestsMicroGinHost -- src/gin/gin_host.cc (GIN_PROXY_NTHREADS). Own
# rccl-UnitTestsMicroGinHost -- src/gin/gin_host.cc (GIN_PROXY_NTHREADS, proxy affinity). Own
# binary: gin-plugin-init-test.cc already defines ncclParamGinEnable in
# rccl-UnitTestsMicro, and gin_fakes.cc supplies ncclGinQueryLastError to Init.
set(TEST_MICRO_GIN_HOST_SOURCE_FILES
Expand Down Expand Up @@ -805,7 +805,7 @@ rccl_add_micro_binary(rccl-UnitTestsMicro
CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml")

# ---------------------------------------------------------------------------
# rccl-UnitTestsMicroGinHost -- src/gin/gin_host.cc (GIN_PROXY_NTHREADS)
# rccl-UnitTestsMicroGinHost -- src/gin/gin_host.cc (GIN_PROXY_NTHREADS, proxy affinity)
# ---------------------------------------------------------------------------
rccl_add_micro_binary(rccl-UnitTestsMicroGinHost
SOURCES ${TEST_MICRO_GIN_HOST_SOURCE_FILES}
Expand Down
30 changes: 29 additions & 1 deletion projects/rccl/test/host/MICROTEST_README.md
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,30 @@ This document is the standing record of:

If you just want to *run* it, jump to [Running and rebuilding](#running-and-rebuilding).

## Naming a new microtest

One unit under test fixes every other name, mechanically, so nothing has to be
invented per binary. For a unit at `src/<dirs>/<stem>.cc`:

| Thing | Rule | `src/gin/gin_host.cc` |
|---|---|---|
| Path macro | `<STEM>_CC_PATH`, `<STEM>` upper-snake | `GIN_HOST_CC_PATH` |
| Test TU | `test/host/<stem>-test.cc`, `_` → `-` | `test/host/gin-host-test.cc` |
| Binary | `rccl-UnitTestsMicro<Stem>`, `<Stem>` UpperCamel | `rccl-UnitTestsMicroGinHost` |
| gtest suites | `<Stem>…Microtest` — every suite in the TU starts with `<Stem>` | `GinHostProxyAffinityMicrotest` |
| CTest categories | `test/test_categories_micro_<stem>.yaml`, lowercase | `test/test_categories_micro_gin_host.yaml` |
| JUnit XML | `host_tests_micro_<stem>.xml` (in `run_host_tests.sh`) | `host_tests_micro_gin_host.xml` |
| Test-runner config | `unit_tests_micro_<stem>` (in `configs/ci-precheckin.json`) | `unit_tests_micro_gin_host` |

Because every suite starts with the binary's `<Stem>`, the categories yaml needs
exactly one pattern (`GinHost*`) and a new suite in the same TU is picked up
without editing it — gtest's `*` does not cross the `.`, so a per-suite list
silently drops any suite someone forgets to add.

A new binary is named in four places: `rccl_add_micro_binary()` in
`test/host/CMakeLists.txt`, the `binaries` array in `run_host_tests.sh`, the
`test_configurations` + `test_suites` pair in `configs/ci-precheckin.json`, and
the microtest list in `lib/test_executor.py` that llvm-cov reads as `--object`.

Comment thread
mch marked this conversation as resolved.
Resolved
## Units under test

Expand Down Expand Up @@ -292,11 +316,15 @@ test:
hipified copy, and `#include` it from the test TU *after* the fakes/macro shims
are in scope. A new unit generally warrants its own binary (see
[Units under test](#units-under-test)) so its file-scope state stays isolated.
The unit's path fixes the test file, binary, suite and yaml names — see
[Naming a new microtest](#naming-a-new-microtest).
2. **Register the source.** Add the test `.cc` to the target's source list in
`rccl_define_micro_source_lists()` in `test/host/CMakeLists.txt`
(`TEST_MICRO_SOURCE_FILES` for `rccl-UnitTestsMicro`), which both build paths
share. If you add a new gtest suite, add its pattern to the target's
`test/test_categories_micro*.yaml` so CTest runs it.
`test/test_categories_micro*.yaml` so CTest runs it (a suite named for its
binary's stem is already matched). A new *binary* also has to be named in the
four places [Naming a new microtest](#naming-a-new-microtest) lists.
3. **Write the `TEST` / fixture.** Use a fixture whose `TearDown()` calls the
unit's reset entry point (`ResetP2pFakes()`, `ResetInitFakes()`, ...) so
hooks do not leak between tests. Install per-test behaviour by overwriting a
Expand Down
167 changes: 163 additions & 4 deletions projects/rccl/test/host/gin-host-test.cc
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,9 @@

// Host-only microtests for src/gin/gin_host.cc GIN_PROXY_NTHREADS (NVIDIA/nccl#2279,
// AICOMRCCL-2017): per-thread endpoint assignment, ginCommCount bump, writer-priority
// list lock, and progress-thread lifecycle.
// list lock, and progress-thread lifecycle, plus proxy-thread CPU affinity
// (AICOMRCCL-1859): the comm->cpuAffinity stash in ginDevCommSetupWithBackend and
// the pin applied by each progress thread.
//
// Own binary (rccl-UnitTestsMicroGinHost): gin-plugin-init-test.cc already defines
// ncclParamGinEnable in rccl-UnitTestsMicro, and gin_fakes.cc supplies
Expand All @@ -16,12 +18,15 @@

#include <gtest/gtest.h>
Comment thread
mch marked this conversation as resolved.
Resolved

#include <sched.h>

#include <atomic>
#include <chrono>
#include <cstdint>
#include <cstdlib>
#include <cstring>
#include <memory>
#include <mutex>
#include <new>
#include <string>
#include <thread>
Expand Down Expand Up @@ -86,12 +91,32 @@ int ncclTeamRankToWorld(ncclComm_t comm, ncclTeam_t team, int rank) {

// Counts entries into ncclGinProgress. The progress loop itself has no counter
// while writePending is set, so tests use this to prove a worker started.
// ncclOsCpuCount is the first thing every progress thread calls, so it doubles as
// the entry counter and as the affinity seam: returning 0 (the default) makes the
// thread skip the pin, non-zero makes it apply ginState->cpuAffinity.
std::atomic<int> g_progressEntries{0};
std::atomic<int> g_osCpuCountValue{0};
// Every mask handed to ncclOsSetAffinity, in call order. Written from the spawned
// progress threads, so it is mutex-guarded; read only after those threads are joined.
std::mutex g_osAffinityMutex;
std::vector<ncclAffinity> g_osSetAffinityMasks;

void ResetAffinityFakes() {
g_progressEntries.store(0);
g_osCpuCountValue.store(0);
std::lock_guard<std::mutex> lock(g_osAffinityMutex);
g_osSetAffinityMasks.clear();
}

int ncclOsCpuCount(const ncclAffinity&) {
g_progressEntries.fetch_add(1, std::memory_order_relaxed);
return 0;
return g_osCpuCountValue.load(std::memory_order_relaxed);
}
ncclResult_t ncclOsSetAffinity(const ncclAffinity& affinity) {
std::lock_guard<std::mutex> lock(g_osAffinityMutex);
g_osSetAffinityMasks.push_back(affinity);
return ncclSuccess;
}
ncclResult_t ncclOsSetAffinity(const ncclAffinity&) { return ncclSuccess; }
void ncclSetThreadName(std::thread&, const char*, ...) {}

#include "fakes/param_redirect.h"
Expand Down Expand Up @@ -237,7 +262,7 @@ class GinHostTest : public ::testing::Test {
g_nLocalGinDevs = 1;
g_peerGinCommCount = -1;
g_paramGinType = -1;
g_progressEntries.store(0);
ResetAffinityFakes();
fake_.holdProgress.store(0);
fake_.progressHolders.store(0);
nthreadsParam_ = 1;
Expand Down Expand Up @@ -634,4 +659,138 @@ TEST_F(GinHostTest, IdleExtraThreadsNeverCallGinProgress) {
freeProgressList();
}

class GinHostProxyAffinityMicrotest : public ::testing::Test {
protected:
ncclGinState ginState_;

void SetUp() override {
ResetAffinityFakes();
CPU_ZERO(&ginState_.cpuAffinity);
ginState_.proxyThreadStopSignal.store(true); // exit at the top of the loop
ginState_.writePending.store(false);
}

void TearDown() override { ResetAffinityFakes(); }

// Run ncclGinProgress on a fresh thread (so the affinity apply, if the fake
// were real, would land on a throwaway thread rather than the test runner)
// and wait for it to return.
void RunProgressOnce() {
std::thread t([this] { ncclGinProgress(&ginState_, /*threadIdx=*/0); });
t.join();
}
};

TEST_F(GinHostProxyAffinityMicrotest, NonEmptyAffinity_PinsProxyThreadToThatCpuSet) {
CPU_SET(3, &ginState_.cpuAffinity);
g_osCpuCountValue.store(1);

RunProgressOnce();

ASSERT_EQ(1u, g_osSetAffinityMasks.size());
EXPECT_TRUE(CPU_ISSET(3, &g_osSetAffinityMasks[0]));
EXPECT_EQ(1, CPU_COUNT(&g_osSetAffinityMasks[0]));
}

TEST_F(GinHostProxyAffinityMicrotest, EmptyAffinity_LeavesProxyThreadAffinityUnchanged) {
g_osCpuCountValue.store(0);

RunProgressOnce();

EXPECT_EQ(1, g_progressEntries.load());
EXPECT_TRUE(g_osSetAffinityMasks.empty());
}

struct FakeGinBackend {
ncclNetDeviceHandle_t devHandle{};
void* ginCtx = reinterpret_cast<void*>(0x1);
int createContextCalls = 0;
};
FakeGinBackend* g_fakeGinBackend = nullptr;

ncclResult_t FakeCreateContext(void* /*collComm*/, ncclGinConfig_t* /*config*/, void** ginCtx,
ncclNetDeviceHandle_t** devHandle) {
g_fakeGinBackend->createContextCalls++;
g_fakeGinBackend->devHandle.handle = reinterpret_cast<void*>(0x2);
g_fakeGinBackend->devHandle.needsProxyProgress = 1;
*ginCtx = g_fakeGinBackend->ginCtx;
*devHandle = &g_fakeGinBackend->devHandle;
return ncclSuccess;
}

ncclResult_t FakeDestroyContext(void* /*ginCtx*/) { return ncclSuccess; }

// The producer: ginDevCommSetupWithBackend copies comm->cpuAffinity into
// ginState->cpuAffinity (gin_host.cc:393) on the branch that first spawns the
// progress threads, so the consumer above has a populated set to read on a real
// comm. Drive that setup path with a scripted backend and assert the copy.
class GinHostProxyAffinitySetupMicrotest : public ::testing::Test {
protected:
std::unique_ptr<ncclComm> comm_ = std::make_unique<ncclComm>();
std::unique_ptr<ncclSharedResources> sr_ = std::make_unique<ncclSharedResources>();
FakeGinBackend fakeBackend_;
ncclGin_t vtable_{};
ncclDevComm devComm_{};

void SetUp() override {
ResetAffinityFakes();
g_fakeGinBackend = &fakeBackend_;

comm_->sharedRes = sr_.get();
struct ncclGinState& ginState = sr_->ginState;
ginState.ginConnectionType = NCCL_GIN_CONNECTION_FULL; // connectedStride == 1
ginState.proxyNthreads = 1;
ginState.proxyThreadsCreated = false; // so needsStart is true
ginState.proxyThreadStopSignal.store(true); // spawned threads exit immediately
ginState.writePending.store(false);

vtable_.name = "fake-gin";
vtable_.createContext = &FakeCreateContext;
vtable_.destroyContext = &FakeDestroyContext;

struct ncclGinBackendState& backend = ginState.backends[0];
backend.ginType = NCCL_GIN_TYPE_PROXY;
backend.ncclGin = &vtable_;
backend.ginCommCount = 1;
backend.ginComms[0] = reinterpret_cast<void*>(0x10);
}

void TearDown() override {
struct ncclGinState& ginState = sr_->ginState;
for (int t = 0; t < ginState.proxyNthreads; t++) {
if (ginState.thread[t].joinable()) ginState.thread[t].join();
}

if (ginState.devComms != nullptr) {
EXPECT_EQ(ncclSuccess, ncclGinDevCommFree(comm_.get(), &devComm_));
}
g_fakeGinBackend = nullptr;
ResetAffinityFakes();
}
};

TEST_F(GinHostProxyAffinitySetupMicrotest, StashesCommAffinityBeforeSpawningProxyThreads) {
CPU_ZERO(&comm_->cpuAffinity);
CPU_SET(5, &comm_->cpuAffinity); // a distinctive mask to spot the copy
g_osCpuCountValue.store(1); // so the spawned worker takes the pin branch

struct ncclGinState& ginState = sr_->ginState;
ncclDevCommRequirements reqs{};
reqs.ginContextCount = 1;
reqs.ginConnectionType = NCCL_GIN_CONNECTION_NONE; // requestedStride stays 1
reqs.ginTrafficClass = NCCL_CONFIG_UNDEF_INT;

ncclResult_t ret =
ginDevCommSetupWithBackend(comm_.get(), &reqs, &devComm_, /*deviceCodeVersion=*/0, &ginState.backends[0]);

ASSERT_EQ(ncclSuccess, ret);
EXPECT_EQ(1, fakeBackend_.createContextCalls); // the setup path really ran
EXPECT_TRUE(ginState.proxyThreadsCreated); // it took the spawn branch
EXPECT_TRUE(CPU_EQUAL(&comm_->cpuAffinity, &ginState.cpuAffinity)); // comm mask stashed verbatim
Comment thread
mch marked this conversation as resolved.
Resolved

ginState.thread[0].join();
ASSERT_EQ(1u, g_osSetAffinityMasks.size());
EXPECT_TRUE(CPU_EQUAL(&comm_->cpuAffinity, &g_osSetAffinityMasks[0]));
}

} // namespace
7 changes: 4 additions & 3 deletions projects/rccl/test/test_categories_micro_gin_host.yaml
Original file line number Diff line number Diff line change
@@ -1,12 +1,13 @@
# rccl-UnitTestsMicroGinHost — shared/ctest schema.
# Host-only micro-tests for src/gin/gin_host.cc GIN_PROXY_NTHREADS (no GPU, no librccl).
# Host-only micro-tests for src/gin/gin_host.cc -- GIN_PROXY_NTHREADS and proxy-thread
# CPU affinity (no GPU, no librccl).
# RCCL does not support Windows: exclude_windows matches all tests so generated suites are empty on Windows.

test_categories:
quick: &micro_gin_host_quick
description: "Smoke — host-only src/gin/gin_host.cc GIN_PROXY_NTHREADS micro-tests, sub-second."
description: "Smoke — host-only src/gin/gin_host.cc GIN_PROXY_NTHREADS and proxy-affinity micro-tests, sub-second."
test_patterns:
- "GinHostTest.*"
- "GinHost*"
exclude_windows:
- "*"
labels:
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -226,7 +226,7 @@
"num_gpus": 1,
"timeout": 60,
"tests": [
{"name": "GinHostMicrotests", "description": "Host-only gin_host.cc GIN_PROXY_NTHREADS microtests (no GPU, no HIP/RCCL runtime); NVIDIA/nccl#2279 / AICOMRCCL-2017", "test_filter": "*"}
{"name": "GinHostMicrotests", "description": "Host-only gin/gin_host.cc microtests (no GPU, no HIP/RCCL runtime); covers GIN_PROXY_NTHREADS (NVIDIA/nccl#2279 / AICOMRCCL-2017) and proxy-thread CPU affinity (AICOMRCCL-1859)", "test_filter": "*"}
]
}
},
Expand Down Expand Up @@ -287,7 +287,7 @@
},
{
"name": "RCCL Microtests - gin_host.cc",
"description": "rccl-UnitTestsMicroGinHost (host-only gin_host.cc GIN_PROXY_NTHREADS microtests, no GPU/HIP/RCCL runtime)",
"description": "rccl-UnitTestsMicroGinHost (host-only gin/gin_host.cc microtests, no GPU/HIP/RCCL runtime)",
"config": "unit_tests_micro_gin_host",
"enabled": true
}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -3291,7 +3291,8 @@ def generate_coverage_report(self):
"rccl-UnitTestsMicroInit-uncached",
"rccl-UnitTestsMicroInit-faultinj",
"rccl-UnitTestsMicroEnqueue",
"rccl-UnitTestsMicroSymKernels"]:
"rccl-UnitTestsMicroSymKernels",
"rccl-UnitTestsMicroGinHost"]:
binary_path = os.path.join(host_test_dir, binary)
if os.path.isfile(binary_path):
object_files.extend(["--object", binary_path])
Expand Down
Loading