Skip to content
Closed
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
34 changes: 34 additions & 0 deletions projects/rccl/test/host/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -293,6 +293,30 @@ function(rccl_define_micro_source_lists)
../common/main_altrsmi.cpp
PARENT_SCOPE)

# rccl-UnitTestsMicroCore -- src/nccl_device/core.cc. Own binary:
# dev_runtime_micro_fakes.cc and devcomm_fakes.cc fake five of its seven functions.
set(TEST_MICRO_CORE_SOURCE_FILES
core-test.cc
fakes/bootstrap_stubs.cc
fakes/topo_stubs.cc
fakes/transport_stubs.cc
fakes/nccl_stubs.cc
fakes/nccl_fakes.cc
fakes/hip_fakes.cc
fakes/collectives_fakes.cc
fakes/comm_fakes.cc
fakes/dev_runtime_fakes.cc
fakes/env_fakes.cc
fakes/os_fakes.cc
fakes/recorder_fakes.cc
fakes/strongstream_stubs.cc
fakes/tuning_fakes.cc
${RCCL_HOST_HIPIFY_DIR}/src/misc/argcheck.cc
${RCCL_HOST_HIPIFY_DIR}/src/misc/archinfo.cc
${RCCL_HOST_HIPIFY_DIR}/src/misc/utils.cc
../common/main_altrsmi.cpp
PARENT_SCOPE)

# rccl-UnitTestsMicroSymKernels -- REAL src/sym_kernels.cc + generated host TU.
set(TEST_MICRO_SYM_KERNELS_SOURCE_FILES
sym-kernels-test.cc
Expand Down Expand Up @@ -1055,6 +1079,16 @@ rccl_add_micro_binary(rccl-UnitTestsMicroClientSupport
CLIENT_SUPPORT_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/ras/client_support.cc"
CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml")

# ---------------------------------------------------------------------------
# rccl-UnitTestsMicroCore
# ---------------------------------------------------------------------------
rccl_add_micro_binary(rccl-UnitTestsMicroCore
SOURCES ${TEST_MICRO_CORE_SOURCE_FILES}
EXTRA_DEFS
${_MICRO_MODE_DEFS}
CORE_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/nccl_device/core.cc"
CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_core.yaml")

# ---------------------------------------------------------------------------
# rccl-UnitTestsMicroSymKernels
# ---------------------------------------------------------------------------
Expand Down
7 changes: 7 additions & 0 deletions projects/rccl/test/host/MICROTEST_README.md
Original file line number Diff line number Diff line change
Expand Up @@ -248,6 +248,11 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary:
guard-less positional pair `libc_seam.h` / `libc_seam_undef.h` must bracket
the production source include, after system headers and before test
helpers. CMake pins the copied `ncclFuncStr` table to `src/init.cc`.
- **`rccl-UnitTestsMicroCore`**: `src/nccl_device/core.cc` (via
`CORE_CC_PATH`, from `core-test.cc`); suite `CoreMicrotest.*`. Its
own binary because `fakes/dev_runtime_micro_fakes.cc` and
`fakes/devcomm_fakes.cc` fake five of its seven functions. See
`test_categories_micro_core.yaml`.
- **`rccl-UnitTestsMicroSymKernels`** — the REAL `src/sym_kernels.cc` (via
`SYM_KERNELS_CC_PATH`, from `sym-kernels-test.cc`), compiled together with the
GENERATED `sym_kernels_host.cc` it calls into; suites `SymKernelMicrotest.*`,
Expand Down Expand Up @@ -829,6 +834,7 @@ can be configured **directly** to build every host binary — `rccl-HostUnitTest
`rccl-UnitTestsMicro`, `rccl-UnitTestsMicroWarpSpeed`,
`rccl-UnitTestsMicroInit[-uncached|-faultinj]`, `rccl-UnitTestsMicroEnqueue[-devlinker]`,
`rccl-UnitTestsMicroGinHost`, `rccl-UnitTestsMicroCe`, `rccl-UnitTestsMicroDiagnostics`,
`rccl-UnitTestsMicroCore`,
`rccl-UnitTestsMicroSymKernels` and `rccl-UnitTestsMicroTaskPrep` — **without configuring/building all of
librccl**. It compiles just the tests + fakes + the hipified unit-under-test
sources.
Expand Down Expand Up @@ -869,6 +875,7 @@ cmake --build build -j"$(nproc)"
./build/rccl-UnitTestsMicroTaskPrep # src/enqueue/task_prep/ + task_sched/ tests
./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/{p2p,ib_write_bw}.cc tests
./build/rccl-UnitTestsMicroGinHost # src/gin/gin_host.cc GIN_PROXY_NTHREADS
./build/rccl-UnitTestsMicroCore # nccl_device/core.cc tests
./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/p2p.cc tests
./build/rccl-UnitTestsMicroCe # src/ce_coll.cc tests
./build/rccl-HostUnitTests
Expand Down
197 changes: 197 additions & 0 deletions projects/rccl/test/host/core-test.cc
Original file line number Diff line number Diff line change
@@ -0,0 +1,197 @@
/*************************************************************************
* Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/

// Host-only microtests for src/nccl_device/core.cc's ncclTeam* and rank-translation helpers.
// ncclDevrInitOnce is faked so both its success and failure paths run.

#include <gtest/gtest.h>

#include <memory>

#include "fakes/dev_runtime_fakes.h"

#include CORE_CC_PATH

namespace {

class CoreMicrotest : public ::testing::Test {
protected:
std::unique_ptr<ncclComm> commStorage;
ncclComm* comm = nullptr;

void SetUp() override {
commStorage = std::make_unique<ncclComm>(); // value-initialised: POD members zeroed
comm = commStorage.get();
}

void TearDown() override { ResetDevRuntimeFakes(); }
};

} // namespace

TEST_F(CoreMicrotest, TeamWorld_ReturnsCommNRanksAndRankWithStrideOne) {
comm->nRanks = 4;
comm->rank = 2;
ncclTeam_t team = ncclTeamWorld(comm);
Comment thread
aduforest marked this conversation as resolved.
Resolved
EXPECT_EQ(team.nRanks, 4);
EXPECT_EQ(team.rank, 2);
EXPECT_EQ(team.stride, 1);
}

TEST_F(CoreMicrotest, TeamWorldAndRankToWorld_IgnoreAFailingInitOnce) {
// The two that don't gate on ncclDevrInitOnce still return real values when it fails.
g_devrInitOnce = [](ncclComm*) { return ncclSystemError; };
comm->nRanks = 4;
comm->rank = 2;
ncclTeam_t world = ncclTeamWorld(comm);
EXPECT_EQ(world.nRanks, 4);
EXPECT_EQ(world.rank, 2);
EXPECT_EQ(world.stride, 1);
ncclTeam_t team{/*nRanks=*/4, /*rank=*/2, /*stride=*/3};
EXPECT_EQ(ncclTeamRankToWorld(comm, team, team.rank + 1), 5); // 2 + (3 - 2) * 3
}

TEST_F(CoreMicrotest, TeamLsa_SuccessPath_ReadsLsaSizeAndSelfFromDevrState) {
comm->devrState.lsaSize = 8;
comm->devrState.lsaSelf = 3;
ncclTeam_t team = ncclTeamLsa(comm);
EXPECT_EQ(team.nRanks, 8);
EXPECT_EQ(team.rank, 3);
EXPECT_EQ(team.stride, 1);
}

TEST_F(CoreMicrotest, TeamLsa_InitOnceFails_ReturnsZeroedTeamWithoutReadingDevrState) {
// Poisoned, so reading them despite the failure shows up below.
comm->devrState.lsaSize = 99;
comm->devrState.lsaSelf = 99;
g_devrInitOnce = [](ncclComm*) { return ncclSystemError; };
ncclTeam_t team = ncclTeamLsa(comm);
EXPECT_EQ(team.nRanks, 0);
EXPECT_EQ(team.rank, 0);
EXPECT_EQ(team.stride, 0);
}

TEST_F(CoreMicrotest, TeamCftMultimem_SuccessPath_ReadsCftMcSizeAndSelfFromDevrState) {
comm->devrState.cftMcSize = 16;
comm->devrState.cftMcSelf = 5;
ncclTeam_t team = ncclTeamCftMultimem(comm);
EXPECT_EQ(team.nRanks, 16);
EXPECT_EQ(team.rank, 5);
EXPECT_EQ(team.stride, 1);
}

TEST_F(CoreMicrotest, TeamCftMultimem_InitOnceFails_ReturnsZeroedTeam) {
comm->devrState.cftMcSize = 99;
comm->devrState.cftMcSelf = 99;
g_devrInitOnce = [](ncclComm*) { return ncclSystemError; };
ncclTeam_t team = ncclTeamCftMultimem(comm);
EXPECT_EQ(team.nRanks, 0);
EXPECT_EQ(team.rank, 0);
EXPECT_EQ(team.stride, 0);
}

TEST_F(CoreMicrotest, TeamRail_SuccessPath_DividesCommNRanksAndRankByLsaSize) {
// 24/4 = 6 and 9/4 = 2 differ from the `%` results ncclTeamInnerFactor would give.
comm->nRanks = 24;
comm->rank = 9;
comm->devrState.lsaSize = 4;
ncclTeam_t team = ncclTeamRail(comm);
EXPECT_EQ(team.nRanks, 6); // 24 / 4
EXPECT_EQ(team.rank, 2); // 9 / 4
EXPECT_EQ(team.stride, 4); // == lsaSize
}

TEST_F(CoreMicrotest, TeamRail_InitOnceFails_ReturnsZeroedTeamWithoutDividingByLsaSize) {
// If the failure branch were missing, this would divide by zero instead.
comm->nRanks = 16;
comm->rank = 10;
comm->devrState.lsaSize = 0;
g_devrInitOnce = [](ncclComm*) { return ncclSystemError; };
ncclTeam_t team = ncclTeamRail(comm);
EXPECT_EQ(team.nRanks, 0);
EXPECT_EQ(team.rank, 0);
EXPECT_EQ(team.stride, 0);
}

TEST_F(CoreMicrotest, TeamCft_FlatMode_ReturnsFlatTeamFromCftSizeAndSelf) {
comm->devrState.cftSize = 12;
comm->devrState.cftSelf = 7;
ncclTeam_t team = ncclTeamCft(comm, NCCL_CFT_TEAM_FLAT);
EXPECT_EQ(team.nRanks, 12);
EXPECT_EQ(team.rank, 7);
EXPECT_EQ(team.stride, 1);
}

TEST_F(CoreMicrotest, TeamCft_HierMultimemMode_OuterFactorsTheFlatTeamByCftMcSize) {
// lsaSize is a distinct non-zero value, so swapped mode arms give a wrong answer, not a crash.
comm->devrState.cftSize = 12;
comm->devrState.cftSelf = 7;
comm->devrState.cftMcSize = 3;
Comment thread
aduforest marked this conversation as resolved.
Resolved
comm->devrState.lsaSize = 2;
ncclTeam_t team = ncclTeamCft(comm, NCCL_CFT_TEAM_HIER_MULTIMEM);
EXPECT_EQ(team.nRanks, 4); // 12 / 3
EXPECT_EQ(team.rank, 2); // 7 / 3
EXPECT_EQ(team.stride, 3); // 1 * 3
}

TEST_F(CoreMicrotest, TeamCft_HierLsaMode_OuterFactorsTheFlatTeamByLsaSize) {
// cftMcSize is distinct for the same reason.
comm->devrState.cftSize = 20;
comm->devrState.cftSelf = 9;
comm->devrState.lsaSize = 5;
comm->devrState.cftMcSize = 3;
ncclTeam_t team = ncclTeamCft(comm, NCCL_CFT_TEAM_HIER_LSA);
EXPECT_EQ(team.nRanks, 4); // 20 / 5
EXPECT_EQ(team.rank, 1); // 9 / 5
EXPECT_EQ(team.stride, 5); // 1 * 5
}

TEST_F(CoreMicrotest, TeamCft_UnrecognisedMode_ReturnsZeroedTeam) {
comm->devrState.cftSize = 12;
comm->devrState.cftSelf = 7;
ncclTeam_t team = ncclTeamCft(comm, static_cast<ncclCftTeamMode_t>(99));
EXPECT_EQ(team.nRanks, 0);
EXPECT_EQ(team.rank, 0);
EXPECT_EQ(team.stride, 0);
}

TEST_F(CoreMicrotest, TeamCft_InitOnceFails_ReturnsZeroedTeamRegardlessOfMode) {
comm->devrState.cftSize = 99;
comm->devrState.cftSelf = 99;
g_devrInitOnce = [](ncclComm*) { return ncclSystemError; };
for (ncclCftTeamMode_t mode : {NCCL_CFT_TEAM_FLAT, NCCL_CFT_TEAM_HIER_MULTIMEM, NCCL_CFT_TEAM_HIER_LSA}) {
ncclTeam_t team = ncclTeamCft(comm, mode);
EXPECT_EQ(team.nRanks, 0);
EXPECT_EQ(team.rank, 0);
EXPECT_EQ(team.stride, 0);
}
}

TEST_F(CoreMicrotest, TeamRankToWorld_MapsATeamLocalRankBackToItsWorldRankViaStride) {
// Pure arithmetic: doesn't gate on ncclDevrInitOnce.
comm->rank = 10;
ncclTeam_t team{/*nRanks=*/4, /*rank=*/2, /*stride=*/3};
// comm->rank + (rank - team.rank) * team.stride
EXPECT_EQ(ncclTeamRankToWorld(comm, team, team.rank), 10); // self: offset 0
EXPECT_EQ(ncclTeamRankToWorld(comm, team, team.rank + 1), 13);
EXPECT_EQ(ncclTeamRankToWorld(comm, team, team.rank - 1), 7);
}

TEST_F(CoreMicrotest, TeamRankToLsa_SuccessPath_MapsATeamLocalRankToAnLsaRankViaLsaSelf) {
comm->devrState.lsaSelf = 5;
ncclTeam_t team{/*nRanks=*/4, /*rank=*/2, /*stride=*/3};
// devrState.lsaSelf + (rank - team.rank) * team.stride
EXPECT_EQ(ncclTeamRankToLsa(comm, team, team.rank), 5);
EXPECT_EQ(ncclTeamRankToLsa(comm, team, team.rank + 1), 8);
Comment thread
aduforest marked this conversation as resolved.
Resolved
EXPECT_EQ(ncclTeamRankToLsa(comm, team, team.rank - 1), 2); // 5 + (-1)*3
}

TEST_F(CoreMicrotest, TeamRankToLsa_InitOnceFails_ReturnsNegativeOneWithoutReadingLsaSelf) {
comm->devrState.lsaSelf = 99;
g_devrInitOnce = [](ncclComm*) { return ncclSystemError; };
ncclTeam_t team{/*nRanks=*/4, /*rank=*/2, /*stride=*/3};
EXPECT_EQ(ncclTeamRankToLsa(comm, team, team.rank), -1);
}
1 change: 1 addition & 0 deletions projects/rccl/test/host/run_host_tests.sh
Original file line number Diff line number Diff line change
Expand Up @@ -167,6 +167,7 @@ do_host_tests() {
"rccl-UnitTestsMicroTaskPrep:$SCRIPT_DIR/host_tests_micro_taskprep.xml"
"rccl-UnitTestsMicroCe:$SCRIPT_DIR/host_tests_micro_ce.xml"
"rccl-UnitTestsMicroClientSupport:$SCRIPT_DIR/host_tests_micro_client_support.xml"
"rccl-UnitTestsMicroCore:$SCRIPT_DIR/host_tests_micro_core.xml"
"rccl-UnitTestsMicroGinHost:$SCRIPT_DIR/host_tests_micro_gin_host.xml"
"rccl-UnitTestsMicroDiagnostics:$SCRIPT_DIR/host_tests_micro_diagnostics.xml"
)
Expand Down
46 changes: 46 additions & 0 deletions projects/rccl/test/test_categories_micro_core.yaml
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
# rccl-UnitTestsMicroCore — shared/ctest schema.
# Host-only micro-tests for src/nccl_device/core.cc (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_core_quick
description: "Smoke — host-only nccl_device/core.cc team-helper micro-tests, sub-second."
test_patterns:
- "Core*"
exclude_windows:
- "*"
labels:
- "rccl"
- "quick"
- "smoke"
- "precheck"
- "microtest"
standard:
<<: *micro_core_quick
labels:
- "rccl"
- "standard"
- "regression"
- "microtest"
comprehensive:
<<: *micro_core_quick
labels:
- "rccl"
- "comprehensive"
- "comprehensive_regression"
- "microtest"
full:
<<: *micro_core_quick
labels:
- "rccl"
- "full"
- "full_regression"
- "microtest"

execution_settings:
timeout_multiplier: 1
category_timeouts:
quick: 600
standard: 600
comprehensive: 600
full: 600
Original file line number Diff line number Diff line change
Expand Up @@ -257,6 +257,17 @@
"tests": [
{"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": "*"}
]
},
"unit_tests_micro_core": {
"is_gtest": true,
"binary": "host/rccl-UnitTestsMicroCore",
"num_ranks": 1,
"num_nodes": 1,
"num_gpus": 1,
"timeout": 60,
"tests": [
{"name": "CoreMicrotests", "description": "Host-only nccl_device/core.cc microtests (no GPU, no HIP/RCCL runtime)", "test_filter": "*"}
]
}
},
"test_suites": [
Expand Down Expand Up @@ -319,6 +330,12 @@
"description": "rccl-UnitTestsMicroGinHost (host-only gin/gin_host.cc microtests, no GPU/HIP/RCCL runtime)",
"config": "unit_tests_micro_gin_host",
"enabled": true
},
{
"name": "RCCL Microtests - nccl_device/core.cc",
"description": "rccl-UnitTestsMicroCore (host-only nccl_device/core.cc microtests, no GPU/HIP/RCCL runtime)",
"config": "unit_tests_micro_core",
"enabled": true
}
]
}
Original file line number Diff line number Diff line change
Expand Up @@ -3292,7 +3292,8 @@ def generate_coverage_report(self):
"rccl-UnitTestsMicroInit-faultinj",
"rccl-UnitTestsMicroEnqueue",
"rccl-UnitTestsMicroSymKernels",
"rccl-UnitTestsMicroGinHost"]:
"rccl-UnitTestsMicroGinHost",
"rccl-UnitTestsMicroCore"]:
binary_path = os.path.join(host_test_dir, binary)
if os.path.isfile(binary_path):
object_files.extend(["--object", binary_path])
Expand Down
Loading