From 8b3ac492961e51500676f9ab3900e6ff5fe40d66 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Tue, 13 Oct 2026 15:54:46 -0700 Subject: [PATCH 01/16] test(rccl): add host microtests for misc/rocmwrap.cc Covers the kernel-version probe, the cuMem VMM functional probe and its cleanup, the cuMem capability gates, the out-of-window NCCL_CUMEM_ENABLE / cuMem-host arms, and rocmLibraryInit's ROCr version and DMA-BUF setup, in a new rccl-UnitTestsMicroRocmwrap binary. Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 23 + projects/rccl/test/host/MICROTEST_README.md | 18 + projects/rccl/test/host/rocmwrap-test.cc | 796 ++++++++++++++++++ projects/rccl/test/host/run_host_tests.sh | 1 + .../test/test_categories_micro_rocmwrap.yaml | 47 ++ .../test_runner/configs/ci-precheckin.json | 17 + 6 files changed, 902 insertions(+) create mode 100644 projects/rccl/test/host/rocmwrap-test.cc create mode 100644 projects/rccl/test/test_categories_micro_rocmwrap.yaml diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 1344ae8b3bf2..dd50b642bfe6 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -388,6 +388,19 @@ function(rccl_define_micro_source_lists) fakes/transport_p2p_fakes.cc # src/transport/p2p.cc shareable buffers ../common/main_altrsmi.cpp PARENT_SCOPE) + + # rccl-UnitTestsMicroRocmwrap -- src/misc/rocmwrap.cc. Its own binary: fakes/rocmwrap_fakes.cc + # and fakes/nccl_fakes.cc (linked by most other targets) define rocmLibraryInit, + # ncclCuMemEnable and ncclCuMemHandleType, which this target compiles for real. The test TU + # supplies ncclDebugLog, g_loadParam and the HSA entry points itself; kernel_config.cc and + # archinfo.cc are the real thing, not fakes. + set(TEST_MICRO_ROCMWRAP_SOURCE_FILES + rocmwrap-test.cc + fakes/hip_fakes.cc + ${RCCL_HOST_HIPIFY_DIR}/src/misc/archinfo.cc # real GetGcnArchName/IsArchMatch + ${RCCL_HOST_HIPIFY_DIR}/src/misc/kernel_config.cc # real option matching + ../common/main_altrsmi.cpp + PARENT_SCOPE) endfunction() # Create one host-only micro-test binary and apply all the boilerplate that @@ -1055,6 +1068,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-UnitTestsMicroRocmwrap +# --------------------------------------------------------------------------- +rccl_add_micro_binary(rccl-UnitTestsMicroRocmwrap + SOURCES ${TEST_MICRO_ROCMWRAP_SOURCE_FILES} + EXTRA_DEFS + ${_MICRO_MODE_DEFS} + ROCMWRAP_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/rocmwrap.cc" + CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_rocmwrap.yaml") + # --------------------------------------------------------------------------- # rccl-UnitTestsMicroSymKernels # --------------------------------------------------------------------------- diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index 1ff07e6588f5..3807aff8f610 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -310,6 +310,22 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: `NCCL_GIN_PROXY_NTHREADS` progress-thread assignment. Its own binary, not sharing `rccl-UnitTestsMicro`: `gin-plugin-init-test.cc` already defines `ncclParamGinEnable` there. See `test_categories_micro_gin_host.yaml`. +- **`rccl-UnitTestsMicroRocmwrap`** — `src/misc/rocmwrap.cc` (via + `ROCMWRAP_CC_PATH`, from `rocmwrap-test.cc`); suite `RocmwrapMicrotest.*`. + Covers the kernel-version probe, the cuMem VMM functional probe and its + cleanup, the capability gates behind `ncclIsCuMemSupported` / + `ncclCuMemRuntimeSupported`, and `rocmLibraryInit`'s ROCr version and + DMA-BUF setup. Its own binary because `fakes/rocmwrap_fakes.cc` and + `fakes/nccl_fakes.cc` define `rocmLibraryInit`, `ncclCuMemEnable` and + `ncclCuMemHandleType`. The test TU defines `ncclDebugLog`, `g_loadParam`, + `hsa_system_get_info` and `hsa_init`, and seams `uname`, `fopen`, + `hipDriverGetVersion`, `hipDeviceGet`, the kernel-config reader and the weak + `hsa_amd_portable_export_dmabuf` symbol inside `rocmwrap.cc` only; + `kernel_config.cc` and `archinfo.cc` are compiled for real. The + `NCCL_CUMEM_VERSION_SUPPORTED(HIP_VERSION)` arms of `ncclCuMemEnable` and + `ncclCuMemHostEnable` compile only on ROCm 7.12+ (or the 7.0.2.x backport), + so on 7.1/7.2 the tests pin the out-of-window arm. See + `test_categories_micro_rocmwrap.yaml`. - **`rccl-UnitTestsMicroDiagnostics`**: `src/diagnostics/p2p.cc` (via `DIAG_P2P_CC_PATH`, suite `DiagP2pMicrotest.*`) and `src/diagnostics/ib_write_bw.cc` (via `DIAG_IB_WRITE_BW_CC_PATH`, suite `DiagIbWriteBwMicrotest.*`). Its own binary: it fakes the @@ -829,6 +845,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-UnitTestsMicroRocmwrap`, `rccl-UnitTestsMicroSymKernels` and `rccl-UnitTestsMicroTaskPrep` — **without configuring/building all of librccl**. It compiles just the tests + fakes + the hipified unit-under-test sources. @@ -869,6 +886,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-UnitTestsMicroRocmwrap # src/misc/rocmwrap.cc cuMem + DMA-BUF init ./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/p2p.cc tests ./build/rccl-UnitTestsMicroCe # src/ce_coll.cc tests ./build/rccl-HostUnitTests diff --git a/projects/rccl/test/host/rocmwrap-test.cc b/projects/rccl/test/host/rocmwrap-test.cc new file mode 100644 index 000000000000..bcf5301b22ae --- /dev/null +++ b/projects/rccl/test/host/rocmwrap-test.cc @@ -0,0 +1,796 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/misc/rocmwrap.cc: the kernel-version probe, the +// cuMem VMM functional probe and capability gates, the NCCL_CUMEM_ENABLE / +// cuMem-host decisions, and rocmLibraryInit's one-time ROCr + DMA-BUF setup. +// +// The real HIP entry points are the shared fakes/hip_fakes.cc hooks. Everything +// else rocmwrap.cc reaches outside HIP goes through a seam scoped to its own +// function bodies (see the #define block before ROCMWRAP_CC_PATH): uname, fopen +// (/proc/kallsyms), hipDriverGetVersion and hipDeviceGet (fixed-value stubs in +// hip_fakes.cc), the kernel-config reader, and the weakly linked +// hsa_amd_portable_export_dmabuf entry. hsa_system_get_info and hsa_init are +// defined here. kernel_config.cc and archinfo.cc are compiled for real, so the +// option matching and the gfx-arch parsing are the production code. +// +// ncclCuMemEnable and ncclCuMemHostEnable each have an arm gated at compile +// time on HIP_VERSION reaching the cuMem window (7.12.60540, or the 7.0.2.x +// backport). ROCm 7.1/7.2 (local and CI) are outside it, so those tests pin the +// out-of-window arm; the in-window arms need hipDeviceAttributeHostNumaId, +// which these headers do not declare. + +#include + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "fakes/hip_fakes.h" + +#include "nccl.h" +#include "debug.h" +#include "rocmwrap.h" +#include "archinfo.h" +#include "kernel_config.h" +#include "hsa/hsa.h" +#include "bootstrap.h" + +// --- Logging and params: this binary links no shared nccl fakes. ------------- + +uint32_t ncclDebugLevelMask = NCCL_DEBUG_LEVEL_MASK_UNINITIALIZED; // log everything +uint64_t ncclDebugMask = ~uint64_t(0); + +namespace { +struct LoggedMessage { + ncclDebugLogLevel level; + std::string text; +}; +std::vector g_logged; +std::map g_params; +} // namespace + +void ncclDebugLog(ncclDebugLogLevel level, unsigned long /*flags*/, const char* /*filefunc*/, int /*line*/, + const char* fmt, ...) { + char buf[512]; + va_list ap; + va_start(ap, fmt); + std::vsnprintf(buf, sizeof(buf), fmt, ap); + va_end(ap); + g_logged.push_back({level, buf}); +} + +std::function g_loadParam = [](const char* env, int64_t deftVal) { + auto it = g_params.find(env); + return it == g_params.end() ? deftVal : it->second; +}; + +#include "fakes/param_redirect.h" + +// --- HSA: the three ROCr entry points rocmwrap.cc calls. ---------------------- + +namespace { +std::function g_hsaSystemGetInfo; +int g_hsaInitCalls = 0; +} // namespace + +extern "C" hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void* value) { + return g_hsaSystemGetInfo(attribute, value); +} +extern "C" hsa_status_t hsa_init() { + ++g_hsaInitCalls; + return HSA_STATUS_SUCCESS; +} + +// --- Seams scoped to rocmwrap.cc's own function bodies. ----------------------- +// Every header rocmwrap.cc includes is pulled in above, so their include guards +// keep the renames below from rewriting any declaration. + +namespace { +std::function g_uname; +std::function g_fopen; +std::function g_driverGetVersion; +std::function g_deviceGet; +std::function g_readKernelConfig; +PFN_hsa_amd_portable_export_dmabuf g_exportDmabuf = nullptr; +} // namespace + +static int RocmWrapTestUname(struct utsname* u) { return g_uname(u); } +static FILE* RocmWrapTestFopen(const char* path, const char* mode) { return g_fopen(path, mode); } +static hipError_t RocmWrapTestDriverGetVersion(int* v) { return g_driverGetVersion(v); } +static hipError_t RocmWrapTestDeviceGet(hipDevice_t* dev, int ordinal) { return g_deviceGet(dev, ordinal); } +static bool RocmWrapTestReadKernelConfig(std::string* content, char* pathOut, size_t pathOutLen) { + return g_readKernelConfig(content, pathOut, pathOutLen); +} + +#define uname RocmWrapTestUname +#define fopen RocmWrapTestFopen +#define hipDriverGetVersion RocmWrapTestDriverGetVersion +#define hipDeviceGet RocmWrapTestDeviceGet +#define ncclKernelConfigReadFirstAvailable RocmWrapTestReadKernelConfig +#define hsa_amd_portable_export_dmabuf g_exportDmabuf +#include ROCMWRAP_CC_PATH +#undef hsa_amd_portable_export_dmabuf +#undef ncclKernelConfigReadFirstAvailable +#undef hipDeviceGet +#undef hipDriverGetVersion +#undef fopen +#undef uname + +namespace { + +constexpr int kKernel_6_8 = KERNEL_VERSION_CODE(6, 8); +constexpr int kCuMemDriverVersion = ROCM_VER_7_12_60540; // inside NCCL_CUMEM_VERSION_SUPPORTED +constexpr int kOldDriverVersion = 70002000; // 7.0.2, below the backport window + +hsa_status_t DmaBufExportStub(const void*, size_t, int*, uint64_t*) { return HSA_STATUS_SUCCESS; } + +// Records each VMM call the functional probe makes, in order. +struct VmmLog { + std::vector calls; + hipMemAllocationProp createProp{}; + hipMemAccessDesc accessDesc{}; + size_t createSize = 0; + std::string failAt; // the call that returns an error; empty means all succeed + size_t granularity = 4096; +}; + +class RocmwrapMicrotest : public ::testing::Test { +protected: + void SetUp() override { + const char* saved = std::getenv("CUDA_LAUNCH_BLOCKING"); + hadLaunchBlocking_ = saved != nullptr; + if (saved) savedLaunchBlocking_ = saved; + Reset(); + } + + void TearDown() override { + if (hadLaunchBlocking_) setenv("CUDA_LAUNCH_BLOCKING", savedLaunchBlocking_.c_str(), 1); + else unsetenv("CUDA_LAUNCH_BLOCKING"); + ResetHipFakes(); + } + + // Every seam back to a working gfx1250 platform on kernel 6.8 with a cuMem-capable driver, + // and rocmwrap.cc's file-scope state back to its initial values. Loops call it per case. + void Reset() { + ResetHipFakes(); + g_logged.clear(); + g_params.clear(); + g_hsaInitCalls = 0; + fopenPaths_.clear(); + unsetenv("CUDA_LAUNCH_BLOCKING"); + + SetKernelRelease("6.8.0-generic"); + g_fopen = [](const char*, const char*) -> FILE* { return nullptr; }; + g_driverGetVersion = [](int* v) { + *v = kCuMemDriverVersion; + return hipSuccess; + }; + g_deviceGet = [](hipDevice_t* dev, int ordinal) { + *dev = ordinal; + return hipSuccess; + }; + g_readKernelConfig = [](std::string*, char*, size_t) { return false; }; + g_exportDmabuf = DmaBufExportStub; + g_hsaSystemGetInfo = [](hsa_system_info_t attribute, void* value) { + if (attribute == HSA_SYSTEM_INFO_VERSION_MAJOR) *static_cast(value) = 1; + else if (attribute == HSA_SYSTEM_INFO_VERSION_MINOR) *static_cast(value) = 14; + else *static_cast(value) = true; // 0x204: DMA-BUF supported + return HSA_STATUS_SUCCESS; + }; + SetArch("gfx1250"); + SetVmmAttribute(1); + defaultVmm_ = VmmLog{}; + InstallVmm(defaultVmm_); + + // rocmwrap.cc's file-scope state. + version_major = version_minor = 0; + ncclCudaLaunchBlocking = false; + ncclCuMemSupported = 0; + ncclCumemHostEnable = -1; + initResult = ncclSuccess; + pfn_hsa_amd_portable_export_dmabuf = nullptr; + } + + static void SetKernelRelease(std::string release) { + g_uname = [release](struct utsname* u) { + std::memset(u, 0, sizeof(*u)); + std::snprintf(u->release, sizeof(u->release), "%s", release.c_str()); + return 0; + }; + } + + static void SetArch(std::string arch) { + g_hipGetDeviceProperties = [arch](hipDeviceProp_t* prop, int) { + *prop = hipDeviceProp_t{}; + std::snprintf(prop->gcnArchName, sizeof(prop->gcnArchName), "%s:sramecc+:xnack-", arch.c_str()); + return hipSuccess; + }; + } + + static void SetVmmAttribute(int value) { + g_hipDeviceGetAttribute = [value](int* pi, hipDeviceAttribute_t attr, int) { + if (attr != hipDeviceAttributeVirtualMemoryManagementSupported) return hipErrorInvalidValue; + *pi = value; + return hipSuccess; + }; + } + + // Installs hooks for every call the functional probe makes, recording each into `log`. + static void InstallVmm(VmmLog& log) { + auto step = [&log](const char* name) { + log.calls.push_back(name); + return log.failAt == name ? hipErrorInvalidValue : hipSuccess; + }; + g_hipMemGetAllocationGranularity = [&log, step](size_t* g, const hipMemAllocationProp* prop, + hipMemAllocationGranularity_flags) { + hipError_t e = step("granularity"); + log.createProp = *prop; + *g = log.granularity; + return e; + }; + g_hipMemCreate = [&log, step](hipMemGenericAllocationHandle_t* h, size_t size, const hipMemAllocationProp*, + unsigned long long) { + log.createSize = size; + *h = reinterpret_cast(0x1); + return step("create"); + }; + g_hipMemAddressReserve = [step](void** ptr, size_t, size_t, void*, unsigned long long) { + *ptr = reinterpret_cast(0x10000); + return step("reserve"); + }; + g_hipMemMap = [step](void*, size_t, size_t, hipMemGenericAllocationHandle_t, unsigned long long) { + return step("map"); + }; + g_hipMemSetAccess = [&log, step](void*, size_t, const hipMemAccessDesc* desc, size_t) { + log.accessDesc = *desc; + return step("setAccess"); + }; + g_hipMemGetAddressRange = [step](hipDeviceptr_t* base, size_t* size, hipDeviceptr_t ptr) { + *base = ptr; + *size = 4096; + return step("addressRange"); + }; + g_hipPointerGetAttribute = [step](void*, hipPointer_attribute, hipDeviceptr_t) { return step("pointerAttr"); }; + g_hipMemUnmap = [step](void*, size_t) { return step("unmap"); }; + g_hipMemAddressFree = [step](void*, size_t) { return step("addressFree"); }; + g_hipMemRelease = [step](hipMemGenericAllocationHandle_t) { return step("release"); }; + } + + static bool Logged(ncclDebugLogLevel level, const char* needle) { + for (const auto& m : g_logged) + if (m.level == level && m.text.find(needle) != std::string::npos) return true; + return false; + } + + // A kernel config that has (or lacks) the two options DMA-BUF needs. + static void SetKernelConfig(bool moveNotify, bool p2pdma) { + g_readKernelConfig = [moveNotify, p2pdma](std::string* content, char* pathOut, size_t pathOutLen) { + *content = "CONFIG_FOO=y\n"; + if (moveNotify) *content += "CONFIG_DMABUF_MOVE_NOTIFY=y\n"; + if (p2pdma) *content += "CONFIG_PCI_P2PDMA=y\n"; + std::snprintf(pathOut, pathOutLen, "/boot/config-test"); + return true; + }; + } + + // /proc/kallsyms served from an in-memory buffer. + void SetKallsyms(std::string content) { + kallsyms_ = std::move(content); + g_fopen = [this](const char* path, const char* mode) -> FILE* { + fopenPaths_.push_back(path); + return fmemopen(kallsyms_.data(), kallsyms_.size(), mode); + }; + } + + std::string kallsyms_; + std::vector fopenPaths_; + VmmLog defaultVmm_; + +private: + bool hadLaunchBlocking_ = false; + std::string savedLaunchBlocking_; +}; + +} // namespace + +// --------------------------------------------------------------------------- +// ncclGetKernelVersionCode +// --------------------------------------------------------------------------- + +TEST_F(RocmwrapMicrotest, KernelVersionCode_PacksMajorAndMinorFromTheRelease) { + SetKernelRelease("6.11.4-45-generic"); + EXPECT_EQ(ncclGetKernelVersionCode(), (6 << 16) | (11 << 8)); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "Kernel version 6.11")); +} + +TEST_F(RocmwrapMicrotest, KernelVersionCode_UnameFailureReturnsMinusOne) { + g_uname = [](struct utsname*) { return -1; }; + EXPECT_EQ(ncclGetKernelVersionCode(), -1); +} + +// sscanf matches nothing, so major and minor keep their zero initialisers. +TEST_F(RocmwrapMicrotest, KernelVersionCode_UnparseableReleaseIsZero) { + SetKernelRelease("custom-kernel"); + EXPECT_EQ(ncclGetKernelVersionCode(), 0); +} + +// --------------------------------------------------------------------------- +// ncclCuMemFunctionalProbe +// --------------------------------------------------------------------------- + +TEST_F(RocmwrapMicrotest, FunctionalProbe_AllStepsSucceed_ReturnsOneAndUndoesEveryStep) { + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclCuMemFunctionalProbe(/*dev=*/0, /*devOrdinal=*/3), 1); + std::vector expected = {"granularity", "create", "reserve", "map", "setAccess", + "addressRange", "pointerAttr", "unmap", "addressFree", "release"}; +#if HIP_VERSION >= 71260540 + expected.insert(expected.begin() + 7, "pointerAttr"); // the legacy-IPC capability query +#endif + EXPECT_EQ(log.calls, expected); + // One granularity unit of pinned device memory on the probed ordinal. + EXPECT_EQ(log.createSize, 4096u); + EXPECT_EQ(log.createProp.type, hipMemAllocationTypePinned); + EXPECT_EQ(log.createProp.location.type, hipMemLocationTypeDevice); + EXPECT_EQ(log.createProp.location.id, 3); + EXPECT_EQ(log.createProp.requestedHandleTypes, ncclCuMemHandleType); + EXPECT_EQ(log.accessDesc.location.type, hipMemLocationTypeDevice); + EXPECT_EQ(log.accessDesc.location.id, 3); + EXPECT_EQ(log.accessDesc.flags, hipMemAccessFlagsProtReadWrite); + EXPECT_FALSE(Logged(NCCL_LOG_INFO, "functional probe failed")); +} + +TEST_F(RocmwrapMicrotest, FunctionalProbe_ZeroGranularityFailsBeforeCreating) { + VmmLog log; + log.granularity = 0; + InstallVmm(log); + EXPECT_EQ(ncclCuMemFunctionalProbe(0, 0), 0); + EXPECT_EQ(log.calls, std::vector({"granularity"})); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "cuMem functional probe failed on device 0")); +} + +// Each failing step undoes exactly the steps that succeeded before it, in reverse order. +TEST_F(RocmwrapMicrotest, FunctionalProbe_EachFailureUndoesOnlyTheCompletedSteps) { + struct Case { + const char* failAt; + std::vector calls; + }; + const std::vector cases = { + {"granularity", {"granularity"}}, + {"create", {"granularity", "create"}}, + {"reserve", {"granularity", "create", "reserve", "release"}}, + {"map", {"granularity", "create", "reserve", "map", "addressFree", "release"}}, + {"setAccess", {"granularity", "create", "reserve", "map", "setAccess", "unmap", "addressFree", "release"}}, + {"addressRange", + {"granularity", "create", "reserve", "map", "setAccess", "addressRange", "unmap", "addressFree", "release"}}, + {"pointerAttr", + {"granularity", "create", "reserve", "map", "setAccess", "addressRange", "pointerAttr", "unmap", "addressFree", + "release"}}, + }; + for (const auto& c : cases) { + SCOPED_TRACE(c.failAt); + g_logged.clear(); + VmmLog log; + log.failAt = c.failAt; + InstallVmm(log); + EXPECT_EQ(ncclCuMemFunctionalProbe(0, 1), 0); + EXPECT_EQ(log.calls, c.calls); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "cuMem functional probe failed on device 1")); + } +} + +// A failing cleanup call is only reported; the probe still answers from the steps before it. +TEST_F(RocmwrapMicrotest, FunctionalProbe_FailingCleanupIsReportedButDoesNotChangeTheAnswer) { + VmmLog log; + log.failAt = "unmap"; + InstallVmm(log); + EXPECT_EQ(ncclCuMemFunctionalProbe(0, 0), 1); + EXPECT_EQ(log.calls.back(), "release"); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "HIP failure")); +} + +// --------------------------------------------------------------------------- +// ncclCuMemCapabilityCheck / ncclIsCuMemSupported +// --------------------------------------------------------------------------- + +TEST_F(RocmwrapMicrotest, CapabilityCheck_AllGatesPassAndProbeSucceeds_ReturnsOne) { + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclCuMemCapabilityCheck(/*requireGfx1250ForAutoEnable=*/1), 1); + EXPECT_FALSE(log.calls.empty()); + EXPECT_FALSE(Logged(NCCL_LOG_WARN, "cuMem support requires")); +} + +TEST_F(RocmwrapMicrotest, CapabilityCheck_AutoEnableOutsideGfx1250_ReturnsZeroWithoutProbing) { + SetArch("gfx942"); + VmmLog log; + InstallVmm(log); + bool unameCalled = false; + g_uname = [&unameCalled](struct utsname*) { + unameCalled = true; + return -1; + }; + EXPECT_EQ(ncclCuMemCapabilityCheck(1), 0); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "limited to gfx1250 (detected gfx942)")); + EXPECT_FALSE(unameCalled); + EXPECT_TRUE(log.calls.empty()); +} + +TEST_F(RocmwrapMicrotest, CapabilityCheck_ArchQueryFailureUnderAutoEnable_ReturnsZero) { + g_hipGetDeviceProperties = [](hipDeviceProp_t*, int) { return hipErrorInvalidDevice; }; + EXPECT_EQ(ncclCuMemCapabilityCheck(1), 0); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "detected unknown")); +} + +// NCCL_CUMEM_ENABLE=1's path skips the arch gate entirely. +TEST_F(RocmwrapMicrotest, CapabilityCheck_WithoutTheGfx1250Gate_AcceptsOtherArchs) { + SetArch("gfx942"); + bool archQueried = false; + g_hipGetDeviceProperties = [&archQueried](hipDeviceProp_t* prop, int) { + archQueried = true; + *prop = hipDeviceProp_t{}; + return hipSuccess; + }; + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclCuMemCapabilityCheck(/*requireGfx1250ForAutoEnable=*/0), 1); + EXPECT_FALSE(archQueried); +} + +// Every cheap gate is evaluated (and WARNs) even after an earlier one failed, but the probe is skipped. +TEST_F(RocmwrapMicrotest, CapabilityCheck_EachCheapGateFailureReturnsZeroWithoutProbing) { + struct Case { + const char* name; + std::function arrange; + const char* warn; + }; + const std::vector cases = { + {"kernel 6.7", [] { SetKernelRelease("6.7.12"); }, "requires Linux kernel >= 6.8"}, + {"driver 7.0.2", + [] { + g_driverGetVersion = [](int* v) { + *v = kOldDriverVersion; + return hipSuccess; + }; + }, + "requires HIP_VERSION >= 7.2.0"}, + {"no VMM", [] { SetVmmAttribute(0); }, "requires VMM RDMA support"}, + }; + for (const auto& c : cases) { + SCOPED_TRACE(c.name); + Reset(); + c.arrange(); + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclCuMemCapabilityCheck(1), 0); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, c.warn)); + EXPECT_TRUE(log.calls.empty()); + } +} + +// The kernel boundary itself: 6.8 passes, so the result rests on the probe alone. +TEST_F(RocmwrapMicrotest, CapabilityCheck_KernelExactly6_8Passes) { + SetKernelRelease("6.8"); + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclCuMemCapabilityCheck(1), 1); + EXPECT_FALSE(Logged(NCCL_LOG_WARN, "kernel")); + static_assert(kKernel_6_8 == ((6 << 16) | (8 << 8)), "KERNEL_VERSION_CODE layout"); +} + +TEST_F(RocmwrapMicrotest, CapabilityCheck_BackportDriverWindowIsAccepted) { + g_driverGetVersion = [](int* v) { + *v = ROCM_VER_7_0_2_2; + return hipSuccess; + }; + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclCuMemCapabilityCheck(1), 1); +} + +TEST_F(RocmwrapMicrotest, CapabilityCheck_ProbeFailureReturnsZero) { + VmmLog log; + log.failAt = "map"; + InstallVmm(log); + EXPECT_EQ(ncclCuMemCapabilityCheck(1), 0); + EXPECT_FALSE(log.calls.empty()); +} + +// HIP query failures jump straight to `error`, which returns 0 since ret is no longer ncclSuccess. +TEST_F(RocmwrapMicrotest, CapabilityCheck_HipQueryFailuresReturnZeroWithoutProbing) { + struct Case { + const char* name; + std::function arrange; + }; + const std::vector cases = { + {"hipGetDevice", [] { g_hipGetDevice = [](int*) { return hipErrorNoDevice; }; }}, + {"hipDriverGetVersion", [] { g_driverGetVersion = [](int*) { return hipErrorNotInitialized; }; }}, + {"hipDeviceGet", [] { g_deviceGet = [](hipDevice_t*, int) { return hipErrorInvalidDevice; }; }}, + {"hipDeviceGetAttribute", + [] { g_hipDeviceGetAttribute = [](int*, hipDeviceAttribute_t, int) { return hipErrorInvalidValue; }; }}, + }; + for (const auto& c : cases) { + SCOPED_TRACE(c.name); + Reset(); + c.arrange(); + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclCuMemCapabilityCheck(0), 0); + EXPECT_TRUE(log.calls.empty()); + } +} + +TEST_F(RocmwrapMicrotest, IsCuMemSupported_AppliesTheGfx1250Gate) { + VmmLog log; + InstallVmm(log); + EXPECT_EQ(ncclIsCuMemSupported(), 1); + SetArch("gfx950"); + EXPECT_EQ(ncclIsCuMemSupported(), 0); +} + +// The memo is process-wide (a function-local once_flag), so check it in a forked child. +TEST_F(RocmwrapMicrotest, RuntimeSupported_SkipsTheArchGateAndRunsTheCheckOnce) { + EXPECT_EXIT( + { + SetArch("gfx942"); + VmmLog log; + InstallVmm(log); + int first = ncclCuMemRuntimeSupported(); + size_t probeCalls = log.calls.size(); + log.failAt = "create"; // would change the answer if the check ran again + int second = ncclCuMemRuntimeSupported(); + bool ok = first == 1 && second == 1 && probeCalls > 0 && log.calls.size() == probeCalls; + std::exit(ok ? 0 : 1); + }, + ::testing::ExitedWithCode(0), ""); +} + +// --------------------------------------------------------------------------- +// ncclCuMemEnable / ncclCuMemHostEnable (out-of-window HIP_VERSION arm) +// --------------------------------------------------------------------------- + +#if !NCCL_CUMEM_VERSION_SUPPORTED(HIP_VERSION) +TEST_F(RocmwrapMicrotest, CuMemEnable_OutsideTheCompileTimeWindowIsAlwaysZero) { + for (int64_t param : {int64_t{-2}, int64_t{0}}) { + g_params["CUMEM_ENABLE"] = param; + g_logged.clear(); + EXPECT_EQ(ncclCuMemEnable(), 0) << param; + EXPECT_FALSE(Logged(NCCL_LOG_WARN, "NCCL_CUMEM_ENABLE=1")) << param; + } + g_params["CUMEM_ENABLE"] = 1; + EXPECT_EQ(ncclCuMemEnable(), 0); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "NCCL_CUMEM_ENABLE=1 is set but cuMem VMM APIs are unavailable")); +} +#endif + +#if !NCCL_CUMEM_HOST_VERSION_SUPPORTED(HIP_VERSION) +TEST_F(RocmwrapMicrotest, CuMemHostEnable_OutsideTheCompileTimeWindowCachesZero) { + EXPECT_EQ(ncclCuMemHostEnable(), 0); + EXPECT_EQ(ncclCumemHostEnable, 0); +} +#endif + +// A cached answer is returned as-is, without re-deciding. +TEST_F(RocmwrapMicrotest, CuMemHostEnable_ReturnsTheCachedValue) { + ncclCumemHostEnable = 1; + EXPECT_EQ(ncclCuMemHostEnable(), 1); +} + +// --------------------------------------------------------------------------- +// initOnceFunc / rocmLibraryInit +// --------------------------------------------------------------------------- + +// CUDA_LAUNCH_BLOCKING is on for any non-empty value other than exactly "0". +TEST_F(RocmwrapMicrotest, Init_ParsesCudaLaunchBlocking) { + struct Case { + const char* value; // nullptr: unset + bool expected; + }; + for (const Case& c : {Case{nullptr, false}, Case{"", false}, Case{"0", false}, Case{"1", true}, Case{"00", true}, + Case{"10", true}}) { + SCOPED_TRACE(c.value ? c.value : ""); + if (c.value) setenv("CUDA_LAUNCH_BLOCKING", c.value, 1); + else unsetenv("CUDA_LAUNCH_BLOCKING"); + ncclCudaLaunchBlocking = !c.expected; + initOnceFunc(); + EXPECT_EQ(ncclCudaLaunchBlocking, c.expected); + } +} + +TEST_F(RocmwrapMicrotest, Init_ReadsTheRocrVersion) { + SetKernelConfig(true, true); + initOnceFunc(); + EXPECT_EQ(version_major, 1); + EXPECT_EQ(version_minor, 14); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "ROCr version 1.14")); +} + +TEST_F(RocmwrapMicrotest, Init_RocrVersionQueryFailureIsASystemErrorBeforeHsaInit) { + for (hsa_system_info_t failing : {HSA_SYSTEM_INFO_VERSION_MAJOR, HSA_SYSTEM_INFO_VERSION_MINOR}) { + SCOPED_TRACE(failing); + Reset(); + g_hsaSystemGetInfo = [failing](hsa_system_info_t attribute, void* value) { + if (attribute == failing) return HSA_STATUS_ERROR; + *static_cast(value) = 1; + return HSA_STATUS_SUCCESS; + }; + initOnceFunc(); + EXPECT_EQ(initResult, ncclSystemError); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "hsa_system_get_info failed")); + EXPECT_EQ(g_hsaInitCalls, 0); + } +} + +// The gfx1250 auto-enable decision is taken once, here, and stored for ncclCuMemEnable. +TEST_F(RocmwrapMicrotest, Init_RecordsTheCuMemCapability) { + SetKernelConfig(true, true); + VmmLog log; + InstallVmm(log); + initOnceFunc(); + EXPECT_EQ(ncclCuMemSupported, 1); + + SetArch("gfx942"); + initOnceFunc(); + EXPECT_EQ(ncclCuMemSupported, 0); +} + +TEST_F(RocmwrapMicrotest, Init_DmaBufDisabledByParamStopsBeforeHsaInit) { + g_params["DMABUF_ENABLE"] = 0; + initOnceFunc(); + EXPECT_EQ(initResult, ncclSystemError); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "Dmabuf feature disabled")); + EXPECT_EQ(g_hsaInitCalls, 0); + EXPECT_EQ(pfn_hsa_amd_portable_export_dmabuf, nullptr); +} + +TEST_F(RocmwrapMicrotest, Init_RocrWithoutDmaBufSupportStopsBeforeHsaInit) { + struct Case { + const char* name; + hsa_status_t status; + bool supported; + }; + for (const Case& c : + {Case{"query fails", HSA_STATUS_ERROR, true}, Case{"reports false", HSA_STATUS_SUCCESS, false}}) { + SCOPED_TRACE(c.name); + Reset(); + g_hsaSystemGetInfo = [c](hsa_system_info_t attribute, void* value) { + if (attribute == HSA_SYSTEM_INFO_VERSION_MAJOR || attribute == HSA_SYSTEM_INFO_VERSION_MINOR) { + *static_cast(value) = 1; + return HSA_STATUS_SUCCESS; + } + EXPECT_EQ(attribute, static_cast(0x204)); + *static_cast(value) = c.supported; + return c.status; + }; + initOnceFunc(); + EXPECT_EQ(initResult, ncclSystemError); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "does not support dmabuf")); + EXPECT_EQ(g_hsaInitCalls, 0); + } +} + +// ROCr advertises DMA-BUF but its export entry did not resolve: disabled cleanly, gate left unarmed. +TEST_F(RocmwrapMicrotest, Init_MissingExportEntryLeavesTheGateUnarmed) { + g_exportDmabuf = nullptr; + initOnceFunc(); + EXPECT_EQ(initResult, ncclSystemError); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "does not export hsa_amd_portable_export_dmabuf")); + EXPECT_EQ(pfn_hsa_amd_portable_export_dmabuf, nullptr); + EXPECT_EQ(g_hsaInitCalls, 0); +} + +TEST_F(RocmwrapMicrotest, Init_KernelConfigWithBothOptionsArmsTheGateAndInitsHsa) { + SetKernelConfig(true, true); + initOnceFunc(); + EXPECT_EQ(initResult, ncclSuccess); + EXPECT_EQ(pfn_hsa_amd_portable_export_dmabuf, &DmaBufExportStub); + EXPECT_EQ(g_hsaInitCalls, 1); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "CONFIG_DMABUF_MOVE_NOTIFY=y in /boot/config-test")); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "CONFIG_PCI_P2PDMA=y in /boot/config-test")); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "DMA_BUF Support Enabled")); + EXPECT_TRUE(fopenPaths_.empty()); +} + +// Pinned as observed: the pfn gate is armed before the kernel check, so it stays armed on this failure. +TEST_F(RocmwrapMicrotest, Init_KernelConfigMissingEitherOptionIsASystemError) { + for (int missing = 0; missing < 2; ++missing) { + SCOPED_TRACE(missing == 0 ? "no MOVE_NOTIFY" : "no P2PDMA"); + Reset(); + SetKernelConfig(missing != 0, missing != 1); + initOnceFunc(); + EXPECT_EQ(initResult, ncclSystemError); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "Failed due to OS kernel support")); + EXPECT_EQ(g_hsaInitCalls, 0); + EXPECT_EQ(pfn_hsa_amd_portable_export_dmabuf, &DmaBufExportStub); + } +} + +// Either order works: each symbol is matched once, whichever line it appears on. +TEST_F(RocmwrapMicrotest, Init_KallsymsFallbackWithBothSymbolsSucceeds) { + for (const char* symbols : {"ffffffff81000000 T _stext\n" + "ffffffff81234560 T dma_buf_move_notify\n" + "ffffffff81234570 T pci_p2pdma_add_resource\n", + "ffffffff81234570 T pci_p2pdma_add_resource\n" + "ffffffff81234580 T pci_p2pdma_distance\n" + "ffffffff81234560 T dma_buf_move_notify\n"}) { + SCOPED_TRACE(symbols); + Reset(); + SetKallsyms(symbols); + initOnceFunc(); + EXPECT_EQ(fopenPaths_, std::vector({"/proc/kallsyms"})); + EXPECT_EQ(initResult, ncclSuccess); + EXPECT_EQ(g_hsaInitCalls, 1); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "via /proc/kallsyms")); + } +} + +TEST_F(RocmwrapMicrotest, Init_KallsymsFallbackMissingEitherSymbolIsASystemError) { + for (const char* symbols : + {"ffffffff81234560 T dma_buf_move_notify\n", "ffffffff81234570 T pci_p2pdma_add_resource\n"}) { + SCOPED_TRACE(symbols); + Reset(); + SetKallsyms(symbols); + initOnceFunc(); + EXPECT_EQ(initResult, ncclSystemError); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "missing kernel symbols in /proc/kallsyms")); + EXPECT_EQ(g_hsaInitCalls, 0); + } +} + +// Pinned as observed: an unreadable /proc/kallsyms is only logged, and init still succeeds. +TEST_F(RocmwrapMicrotest, Init_UnreadableKallsymsStillSucceeds) { + initOnceFunc(); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "Could not open /proc/kallsyms")); + EXPECT_EQ(initResult, ncclSuccess); + EXPECT_EQ(g_hsaInitCalls, 1); +} + +// RCCL_FORCE_ENABLE_DMABUF=1 overrides NCCL_DMABUF_ENABLE=0 and skips the kernel checks. +TEST_F(RocmwrapMicrotest, Init_ForceEnableSkipsTheParamAndKernelChecks) { + g_params["RCCL_FORCE_ENABLE_DMABUF"] = 1; + g_params["DMABUF_ENABLE"] = 0; + bool configRead = false; + g_readKernelConfig = [&configRead](std::string*, char*, size_t) { + configRead = true; + return false; + }; + initOnceFunc(); + EXPECT_EQ(initResult, ncclSuccess); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "DMA_BUF Support is force enabled")); + EXPECT_FALSE(configRead); + EXPECT_TRUE(fopenPaths_.empty()); + EXPECT_EQ(g_hsaInitCalls, 1); +} + +// pthread_once state is process-wide, so check the once-only behaviour in a forked child. +TEST_F(RocmwrapMicrotest, LibraryInit_RunsInitOnceAndReturnsItsResult) { + EXPECT_EXIT( + { + int majorQueries = 0; + g_hsaSystemGetInfo = [&majorQueries](hsa_system_info_t attribute, void* value) { + if (attribute == HSA_SYSTEM_INFO_VERSION_MAJOR) ++majorQueries; + *static_cast(value) = 1; + return attribute == HSA_SYSTEM_INFO_VERSION_MINOR ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS; + }; + ncclResult_t first = rocmLibraryInit(); + ncclResult_t second = rocmLibraryInit(); + bool ok = first == ncclSystemError && second == ncclSystemError && majorQueries == 1; + std::exit(ok ? 0 : 1); + }, + ::testing::ExitedWithCode(0), ""); +} diff --git a/projects/rccl/test/host/run_host_tests.sh b/projects/rccl/test/host/run_host_tests.sh index ec702f4f62b1..be109a8d97de 100755 --- a/projects/rccl/test/host/run_host_tests.sh +++ b/projects/rccl/test/host/run_host_tests.sh @@ -168,6 +168,7 @@ do_host_tests() { "rccl-UnitTestsMicroCe:$SCRIPT_DIR/host_tests_micro_ce.xml" "rccl-UnitTestsMicroClientSupport:$SCRIPT_DIR/host_tests_micro_client_support.xml" "rccl-UnitTestsMicroGinHost:$SCRIPT_DIR/host_tests_micro_gin_host.xml" + "rccl-UnitTestsMicroRocmwrap:$SCRIPT_DIR/host_tests_micro_rocmwrap.xml" "rccl-UnitTestsMicroDiagnostics:$SCRIPT_DIR/host_tests_micro_diagnostics.xml" ) # Binaries that only exist for some CMake option settings (rccl-UnitTestsMicroSymKernels needs diff --git a/projects/rccl/test/test_categories_micro_rocmwrap.yaml b/projects/rccl/test/test_categories_micro_rocmwrap.yaml new file mode 100644 index 000000000000..d89c9fe6d282 --- /dev/null +++ b/projects/rccl/test/test_categories_micro_rocmwrap.yaml @@ -0,0 +1,47 @@ +# rccl-UnitTestsMicroRocmwrap — shared/ctest schema. +# Host-only micro-tests for src/misc/rocmwrap.cc -- kernel-version probe, cuMem VMM capability +# checks, and rocmLibraryInit's ROCr/DMA-BUF setup (no GPU, no librccl). +# RCCL does not support Windows: exclude_windows matches all tests so generated suites are empty on Windows. + +test_categories: + quick: µ_rocmwrap_quick + description: "Smoke — host-only src/misc/rocmwrap.cc cuMem capability and DMA-BUF init micro-tests, sub-second." + test_patterns: + - "Rocmwrap*" + exclude_windows: + - "*" + labels: + - "rccl" + - "quick" + - "smoke" + - "precheck" + - "microtest" + standard: + <<: *micro_rocmwrap_quick + labels: + - "rccl" + - "standard" + - "regression" + - "microtest" + comprehensive: + <<: *micro_rocmwrap_quick + labels: + - "rccl" + - "comprehensive" + - "comprehensive_regression" + - "microtest" + full: + <<: *micro_rocmwrap_quick + labels: + - "rccl" + - "full" + - "full_regression" + - "microtest" + +execution_settings: + timeout_multiplier: 1 + category_timeouts: + quick: 600 + standard: 600 + comprehensive: 600 + full: 600 diff --git a/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json b/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json index 07da6b28ba73..dc09e7f69e11 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json +++ b/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json @@ -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_rocmwrap": { + "is_gtest": true, + "binary": "host/rccl-UnitTestsMicroRocmwrap", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 60, + "tests": [ + {"name": "RocmwrapMicrotests", "description": "Host-only misc/rocmwrap.cc microtests (no GPU, no HIP/RCCL runtime); covers the cuMem VMM capability checks and rocmLibraryInit's ROCr/DMA-BUF setup", "test_filter": "*"} + ] } }, "test_suites": [ @@ -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 - rocmwrap.cc", + "description": "rccl-UnitTestsMicroRocmwrap (host-only misc/rocmwrap.cc microtests, no GPU/HIP/RCCL runtime)", + "config": "unit_tests_micro_rocmwrap", + "enabled": true } ] } From 17147e87066458a8d7745794e2896d502c238c00 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Tue, 13 Oct 2026 19:38:36 -0700 Subject: [PATCH 02/16] test(rccl): mark kernel_config.cc generated for in-tree configure, address ARGUS findings on rocmwrap microtests Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 4 ++-- projects/rccl/test/host/MICROTEST_README.md | 9 +++++---- projects/rccl/test/host/rocmwrap-test.cc | 20 ++++++++++++------- .../scripts/test_runner/lib/test_executor.py | 3 ++- 4 files changed, 22 insertions(+), 14 deletions(-) diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index dd50b642bfe6..88e618ffd12f 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -873,7 +873,7 @@ endif() # Mark every hipified misc/ oracle TU any target below compiles as GENERATED. # The property is directory-scoped and idempotent, so one call up front covers # all targets -- no need to repeat it per binary. -rccl_mark_oracle_tus_generated(${RCCL_HOST_HIPIFY_DIR} argcheck archinfo utils) +rccl_mark_oracle_tus_generated(${RCCL_HOST_HIPIFY_DIR} argcheck archinfo utils kernel_config) # --------------------------------------------------------------------------- # rccl-UnitTestsMicro @@ -1103,7 +1103,7 @@ endif() # in rccl_define_micro_source_lists() for the rationale. # --------------------------------------------------------------------------- # task_prep/ and config/ oracle TUs live outside src/misc/, so mark them GENERATED -# here (rccl_mark_oracle_tus_generated only covers the misc/ argcheck/archinfo/utils set). +# here (rccl_mark_oracle_tus_generated only covers the misc/ argcheck/archinfo/utils/kernel_config set). set_source_files_properties( ${RCCL_HOST_HIPIFY_DIR}/src/config/algorithm_parser.cc ${RCCL_HOST_HIPIFY_DIR}/src/config/algorithm_registry.cc diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index 3807aff8f610..e8bd2dba986c 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -321,10 +321,11 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: `hsa_system_get_info` and `hsa_init`, and seams `uname`, `fopen`, `hipDriverGetVersion`, `hipDeviceGet`, the kernel-config reader and the weak `hsa_amd_portable_export_dmabuf` symbol inside `rocmwrap.cc` only; - `kernel_config.cc` and `archinfo.cc` are compiled for real. The - `NCCL_CUMEM_VERSION_SUPPORTED(HIP_VERSION)` arms of `ncclCuMemEnable` and - `ncclCuMemHostEnable` compile only on ROCm 7.12+ (or the 7.0.2.x backport), - so on 7.1/7.2 the tests pin the out-of-window arm. See + `kernel_config.cc` and `archinfo.cc` are compiled for real. The in-window + arm of `ncclCuMemEnable` compiles only on ROCm 7.12+ or the 7.0.2.x backport + (`NCCL_CUMEM_VERSION_SUPPORTED`), and that of `ncclCuMemHostEnable` only on + 7.12+ (`NCCL_CUMEM_HOST_VERSION_SUPPORTED`), so on 7.1/7.2 the tests pin the + out-of-window arms. See `test_categories_micro_rocmwrap.yaml`. - **`rccl-UnitTestsMicroDiagnostics`**: `src/diagnostics/p2p.cc` (via `DIAG_P2P_CC_PATH`, suite `DiagP2pMicrotest.*`) and `src/diagnostics/ib_write_bw.cc` diff --git a/projects/rccl/test/host/rocmwrap-test.cc b/projects/rccl/test/host/rocmwrap-test.cc index bcf5301b22ae..4ad99b6c8ac8 100644 --- a/projects/rccl/test/host/rocmwrap-test.cc +++ b/projects/rccl/test/host/rocmwrap-test.cc @@ -18,10 +18,11 @@ // option matching and the gfx-arch parsing are the production code. // // ncclCuMemEnable and ncclCuMemHostEnable each have an arm gated at compile -// time on HIP_VERSION reaching the cuMem window (7.12.60540, or the 7.0.2.x -// backport). ROCm 7.1/7.2 (local and CI) are outside it, so those tests pin the -// out-of-window arm; the in-window arms need hipDeviceAttributeHostNumaId, -// which these headers do not declare. +// time on HIP_VERSION: ncclCuMemEnable's on the cuMem window (7.12.60540, or the +// 7.0.2.x backport), ncclCuMemHostEnable's on 7.12.60540 only. ROCm 7.1/7.2 +// (local and CI) are outside both, so those tests pin the out-of-window arms; +// the host arm also needs hipDeviceAttributeHostNumaId, which these headers do +// not declare. #include @@ -175,7 +176,10 @@ class RocmwrapMicrotest : public ::testing::Test { unsetenv("CUDA_LAUNCH_BLOCKING"); SetKernelRelease("6.8.0-generic"); - g_fopen = [](const char*, const char*) -> FILE* { return nullptr; }; + g_fopen = [this](const char* path, const char*) -> FILE* { + fopenPaths_.push_back(path); + return nullptr; + }; g_driverGetVersion = [](int* v) { *v = kCuMemDriverVersion; return hipSuccess; @@ -272,8 +276,9 @@ class RocmwrapMicrotest : public ::testing::Test { } static bool Logged(ncclDebugLogLevel level, const char* needle) { - for (const auto& m : g_logged) + for (const auto& m : g_logged) { if (m.level == level && m.text.find(needle) != std::string::npos) return true; + } return false; } @@ -339,7 +344,7 @@ TEST_F(RocmwrapMicrotest, FunctionalProbe_AllStepsSucceed_ReturnsOneAndUndoesEve EXPECT_EQ(ncclCuMemFunctionalProbe(/*dev=*/0, /*devOrdinal=*/3), 1); std::vector expected = {"granularity", "create", "reserve", "map", "setAccess", "addressRange", "pointerAttr", "unmap", "addressFree", "release"}; -#if HIP_VERSION >= 71260540 +#if HIP_VERSION >= ROCM_VER_7_12_60540 expected.insert(expected.begin() + 7, "pointerAttr"); // the legacy-IPC capability query #endif EXPECT_EQ(log.calls, expected); @@ -444,6 +449,7 @@ TEST_F(RocmwrapMicrotest, CapabilityCheck_WithoutTheGfx1250Gate_AcceptsOtherArch g_hipGetDeviceProperties = [&archQueried](hipDeviceProp_t* prop, int) { archQueried = true; *prop = hipDeviceProp_t{}; + std::snprintf(prop->gcnArchName, sizeof(prop->gcnArchName), "gfx942:sramecc+:xnack-"); return hipSuccess; }; VmmLog log; diff --git a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py index 70add7053ab2..2b85e2f1ed93 100644 --- a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py +++ b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py @@ -3292,7 +3292,8 @@ def generate_coverage_report(self): "rccl-UnitTestsMicroInit-faultinj", "rccl-UnitTestsMicroEnqueue", "rccl-UnitTestsMicroSymKernels", - "rccl-UnitTestsMicroGinHost"]: + "rccl-UnitTestsMicroGinHost", + "rccl-UnitTestsMicroRocmwrap"]: binary_path = os.path.join(host_test_dir, binary) if os.path.isfile(binary_path): object_files.extend(["--object", binary_path]) From 8559f6a70d9e39d7cd2369e2ba41d724b6930c31 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Tue, 13 Oct 2026 00:53:31 -0700 Subject: [PATCH 03/16] test(rccl): add host microtests for misc/roctx.cc and misc/nvmlwrap_stub.cc New rccl-UnitTestsMicroRoctxNvml binary compiling both files for real (comm_fakes.cc stubs roctx_scoped_range_in for every other target). The test TU supplies its own ncclDebugLog and g_loadParam, so no shared fakes are linked; param_redirect.h makes RCCL_LOG_ROCTX controllable per test, and a malloc seam scoped to roctx.cc reaches both roctxAlloc failure paths. stringify() overflows its 1024-byte buffer once a message is truncated (confirmed under ASan); filed as AICOMRCCL-2914 and not pinned here. Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 21 + .../rccl/test/host/roctx-nvml-stub-test.cc | 528 ++++++++++++++++++ projects/rccl/test/host/run_host_tests.sh | 1 + projects/rccl/test/test_categories_micro.yaml | 3 + 4 files changed, 553 insertions(+) create mode 100644 projects/rccl/test/host/roctx-nvml-stub-test.cc diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 1344ae8b3bf2..9621438ef4cf 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -293,6 +293,16 @@ function(rccl_define_micro_source_lists) ../common/main_altrsmi.cpp PARENT_SCOPE) + # rccl-UnitTestsMicroRoctxNvml -- src/misc/roctx.cc + src/misc/nvmlwrap_stub.cc. Its own binary: + # fakes/comm_fakes.cc (linked by rccl-UnitTestsMicro and most other targets) stubs the + # roctx_scoped_range_in constructor/destructor this target compiles as the real thing, and + # ras-diagnostics-gpu-test.cc defines its own NVML seams. The test TU supplies ncclDebugLog and + # g_loadParam itself, so no shared fakes are linked. + set(TEST_MICRO_ROCTX_NVML_SOURCE_FILES + roctx-nvml-stub-test.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 @@ -1055,6 +1065,17 @@ 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-UnitTestsMicroRoctxNvml +# --------------------------------------------------------------------------- +rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml + SOURCES ${TEST_MICRO_ROCTX_NVML_SOURCE_FILES} + EXTRA_DEFS + ${_MICRO_MODE_DEFS} + ROCTX_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/roctx.cc" + NVMLWRAP_STUB_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/nvmlwrap_stub.cc" + CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml") + # --------------------------------------------------------------------------- # rccl-UnitTestsMicroSymKernels # --------------------------------------------------------------------------- diff --git a/projects/rccl/test/host/roctx-nvml-stub-test.cc b/projects/rccl/test/host/roctx-nvml-stub-test.cc new file mode 100644 index 000000000000..6f375cec7e91 --- /dev/null +++ b/projects/rccl/test/host/roctx-nvml-stub-test.cc @@ -0,0 +1,528 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/misc/roctx.cc (NVTX payload -> roctx message +// translation and the RAII range used by NCCL_NVTX3_FUNC_RANGE*) and +// src/misc/nvmlwrap_stub.cc (the ROCm stand-in for NVML, which reports no +// devices so NVML-driven callers skip their checks). +// +// Both files are compiled for real into this binary. Neither needs the shared +// fakes: this TU supplies its own ncclDebugLog (to record WARNs) and +// g_loadParam (param_redirect.h routes rcclParamLogRoctx() through it on every +// call instead of caching, so each test picks LOG_ROCTX independently). +// +// The host build never defines ROCTX_ENABLE, so roctxRangePushA/Pop compile out +// and the scoped range's observable behaviour on this build is its WARN +// fallback plus the payload alloc/free it performs around it. +// +// Every stringify() case here stays under MAX_MESSAGE_LENGTH on purpose: past +// it, stringify() writes beyond its 1024-byte buffer (confirmed under ASan), +// tracked as AICOMRCCL-2914 rather than pinned here as undefined behaviour. + +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "debug.h" + +std::function g_loadParam = [](const char*, int64_t deftVal) { return deftVal; }; + +namespace { +struct LoggedMessage { + ncclDebugLogLevel level; + std::string text; +}; +std::vector g_logged; +} // namespace + +void ncclDebugLog(ncclDebugLogLevel level, unsigned long /*flags*/, const char* /*filefunc*/, int /*line*/, + const char* fmt, ...) { + char buf[512]; + va_list ap; + va_start(ap, fmt); + std::vsnprintf(buf, sizeof(buf), fmt, ap); + va_end(ap); + g_logged.push_back({level, buf}); +} + +#include "fakes/param_redirect.h" + +// malloc seam for roctx.cc only, so roctxAlloc's two failure paths are reachable. Every header +// roctx.cc includes is pulled in first, so their include guards keep the macro below from +// rewriting any declaration; it applies to roctx.cc's own function bodies and nothing else. +#include +#include "roctx.h" +namespace { +std::function g_roctxMalloc; +} // namespace +static void* RoctxTestMalloc(size_t n) { return g_roctxMalloc ? g_roctxMalloc(n) : std::malloc(n); } +#define malloc(n) RoctxTestMalloc(n) +#include ROCTX_CC_PATH +#undef malloc + +#include NVMLWRAP_STUB_CC_PATH + +namespace { + +class RoctxMicrotest : public ::testing::Test { +protected: + void SetUp() override { g_logged.clear(); } + void TearDown() override { + g_loadParam = [](const char*, int64_t deftVal) { return deftVal; }; + g_roctxMalloc = nullptr; + g_logged.clear(); + } + + // Lets the first `succeedFirst` roctx.cc mallocs through, then fails every one after. + static int* FailMallocAfter(int succeedFirst) { + static int calls; + calls = 0; + g_roctxMalloc = [succeedFirst](size_t n) -> void* { return calls++ < succeedFirst ? std::malloc(n) : nullptr; }; + return &calls; + } + + static void SetLogRoctx(int64_t value) { + g_loadParam = [value](const char* env, int64_t deftVal) { + return std::strcmp(env, "RCCL_LOG_ROCTX") == 0 ? value : deftVal; + }; + } + + static size_t WarnCount() { + size_t n = 0; + for (const auto& m : g_logged) n += (m.level == NCCL_LOG_WARN); + return n; + } + + // Zero-initialised schema entry: the fields roctx.cc does not read stay at 0. + static nvtxPayloadSchemaEntry_t Entry(uint64_t type, const char* name, uint64_t offset) { + nvtxPayloadSchemaEntry_t e; + std::memset(&e, 0, sizeof(e)); + e.type = type; + e.name = name; + e.offset = offset; + return e; + } +}; + +struct MixedPayload { + int i; + size_t s; + ncclDevRedOp_t op; + ncclDataType_t dt; +}; + +} // namespace + +// --------------------------------------------------------------------------- +// Lookup tables +// --------------------------------------------------------------------------- + +TEST_F(RoctxMicrotest, NvtxToRoctx_MapsEachSupportedNvtxTypeToItsRoctxType) { + EXPECT_EQ(nvtxToRoctx.size(), 4u); + EXPECT_EQ(nvtxToRoctx.at(NVTX_PAYLOAD_ENTRY_TYPE_INT), ROCTX_PAYLOAD_ENTRY_TYPE_INT); + EXPECT_EQ(nvtxToRoctx.at(NVTX_PAYLOAD_ENTRY_TYPE_SIZE), ROCTX_PAYLOAD_ENTRY_TYPE_SIZE); + EXPECT_EQ(nvtxToRoctx.at(NVTX_PAYLOAD_ENTRY_TYPE_REDOP), ROCTX_PAYLOAD_ENTRY_TYPE_REDOP); + EXPECT_EQ(nvtxToRoctx.at(NVTX_PAYLOAD_ENTRY_TYPE_DATATYPE), ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE); +} + +// The names stringify() prints for redop/datatype payloads, indexed by enum value. +TEST_F(RoctxMicrotest, NameTables_MatchTheEnumOrder) { + const char* redOps[] = {"Sum", "Prod", "MinMax", "PreMulSum", "SumPostDiv"}; + ASSERT_EQ(sizeof(redOps) / sizeof(redOps[0]), static_cast(ncclNumDevRedOps)); + for (int i = 0; i < ncclNumDevRedOps; ++i) EXPECT_STREQ(ncclRedOpStr[i], redOps[i]) << "op " << i; + + const char* dataTypes[] = {"i8", "u8", "i32", "u32", "i64", "u64", "f16", "f32", "f64", "b16", "f8", "b8"}; + ASSERT_EQ(sizeof(dataTypes) / sizeof(dataTypes[0]), static_cast(ncclNumTypes)); + for (int i = 0; i < ncclNumTypes; ++i) EXPECT_STREQ(ncclDataTypeStr[i], dataTypes[i]) << "type " << i; +} + +// Pinned as observed, not endorsed: the table has 4 slots but only 3 initialisers, so the +// DATATYPE slot is null. Nothing in the tree reads roctxEntryTypeStr today. +TEST_F(RoctxMicrotest, RoctxEntryTypeStr_DatatypeSlotIsNull) { + EXPECT_STREQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_INT], "ROCTX_PAYLOAD_ENTRY_TYPE_INT"); + EXPECT_STREQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_SIZE], "ROCTX_PAYLOAD_ENTRY_TYPE_SIZE"); + EXPECT_STREQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_REDOP], "ROCTX_PAYLOAD_ENTRY_TYPE_REDOP"); + EXPECT_EQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE], nullptr); +} + +// --------------------------------------------------------------------------- +// roctxAlloc / roctxFree +// --------------------------------------------------------------------------- + +TEST_F(RoctxMicrotest, Alloc_AllocatesEntriesAndMessage_FreeReleasesBoth) { + roctxPayloadInfo info; + roctxAlloc(&info, 3); + ASSERT_NE(info.payloadEntries, nullptr); + ASSERT_NE(info.message, nullptr); + // Both buffers are usable at their full advertised size (ASan flags a short allocation). + std::memset(info.payloadEntries, 0, 3 * sizeof(roctxPayloadSchemaEntryInfo)); + std::memset(info.message, 'x', MAX_MESSAGE_LENGTH); + roctxFree(&info); +} + +// The entries allocation fails: both pointers end up null and the message is never allocated. +TEST_F(RoctxMicrotest, Alloc_EntriesAllocationFails_LeavesBothPointersNull) { + int* calls = FailMallocAfter(0); + roctxPayloadInfo info; + info.message = reinterpret_cast(0x1); // must be overwritten with nullptr + roctxAlloc(&info, 2); + EXPECT_EQ(info.payloadEntries, nullptr); + EXPECT_EQ(info.message, nullptr); + EXPECT_EQ(*calls, 1); +} + +// The message allocation fails after the entries succeeded: the entries buffer is released +// (LeakSanitizer catches a missing free) and nulled, so a later roctxFree cannot double-free. +TEST_F(RoctxMicrotest, Alloc_MessageAllocationFails_ReleasesEntriesAndNullsBoth) { + int* calls = FailMallocAfter(1); + roctxPayloadInfo info; + roctxAlloc(&info, 2); + EXPECT_EQ(*calls, 2); + EXPECT_EQ(info.payloadEntries, nullptr); + EXPECT_EQ(info.message, nullptr); + roctxFree(&info); +} + +TEST_F(RoctxMicrotest, Free_ToleratesNullPointers) { + roctxPayloadInfo info; // both members default to nullptr + roctxFree(&info); + EXPECT_EQ(info.payloadEntries, nullptr); + EXPECT_EQ(info.message, nullptr); +} + +// --------------------------------------------------------------------------- +// extractPayloadInfo / stringify +// --------------------------------------------------------------------------- + +TEST_F(RoctxMicrotest, Extract_NoEntriesBuffer_ReturnsWithoutTouchingInfo) { + roctxPayloadInfo info; + info.id = "untouched"; + info.numEntries = 77; + nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "a", 0)}; + int value = 1; + nvtxPayloadData_t data{0, sizeof(value), &value}; + extractPayloadInfo(schema, &data, 1, "Sched", &info); + EXPECT_STREQ(info.id, "untouched"); + EXPECT_EQ(info.numEntries, 77u); +} + +TEST_F(RoctxMicrotest, Extract_ReadsEachTypeAtItsOffsetAndStringifies) { + MixedPayload payload{-3, size_t{1} << 40, ncclDevProd, ncclFloat32}; + nvtxPayloadSchemaEntry_t schema[] = { + Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "count", offsetof(MixedPayload, i)), + Entry(NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "bytes", offsetof(MixedPayload, s)), + Entry(NVTX_PAYLOAD_ENTRY_TYPE_REDOP, "op", offsetof(MixedPayload, op)), + Entry(NVTX_PAYLOAD_ENTRY_TYPE_DATATYPE, "type", offsetof(MixedPayload, dt)), + }; + nvtxPayloadData_t data{0, sizeof(payload), &payload}; + roctxPayloadInfo info; + roctxAlloc(&info, 4); + extractPayloadInfo(schema, &data, 4, "AllReduce", &info); + + EXPECT_STREQ(info.id, "AllReduce"); + ASSERT_EQ(info.numEntries, 4u); + EXPECT_STREQ(info.payloadEntries[0].name, "count"); + EXPECT_EQ(info.payloadEntries[0].type, ROCTX_PAYLOAD_ENTRY_TYPE_INT); + EXPECT_EQ(info.payloadEntries[0].payload.typeInt, -3); + EXPECT_EQ(info.payloadEntries[1].type, ROCTX_PAYLOAD_ENTRY_TYPE_SIZE); + EXPECT_EQ(info.payloadEntries[1].payload.typeSize, size_t{1} << 40); + EXPECT_EQ(info.payloadEntries[2].type, ROCTX_PAYLOAD_ENTRY_TYPE_REDOP); + EXPECT_EQ(info.payloadEntries[2].payload.typeRedOp, ncclDevProd); + EXPECT_EQ(info.payloadEntries[3].type, ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE); + EXPECT_EQ(info.payloadEntries[3].payload.typeDataType, ncclFloat32); + EXPECT_STREQ(info.message, "{AllReduce: count: -3, bytes: 1099511627776, op: Prod, type: f32}"); + roctxFree(&info); +} + +// The separator is only written between entries: one entry has no trailing ", ". +TEST_F(RoctxMicrotest, Stringify_SingleEntryHasNoSeparator) { + roctxPayloadInfo info; + roctxAlloc(&info, 1); + info.id = "Bcast"; + info.numEntries = 1; + info.payloadEntries[0].name = "root"; + info.payloadEntries[0].type = ROCTX_PAYLOAD_ENTRY_TYPE_INT; + info.payloadEntries[0].payload.typeInt = 7; + stringify(&info); + EXPECT_STREQ(info.message, "{Bcast: root: 7}"); + roctxFree(&info); +} + +TEST_F(RoctxMicrotest, Stringify_ZeroEntriesIsJustTheId) { + roctxPayloadInfo info; + roctxAlloc(&info, 1); + info.id = "Empty"; + info.numEntries = 0; + stringify(&info); + EXPECT_STREQ(info.message, "{Empty: }"); + roctxFree(&info); +} + +// Out-of-range redop/datatype values print "unknown" instead of indexing past the tables; +// the highest valid values sit right next to them so an off-by-one in either bound shows. +TEST_F(RoctxMicrotest, Stringify_OutOfRangeRedOpAndDataTypePrintUnknown) { + roctxPayloadInfo info; + roctxAlloc(&info, 4); + info.id = "X"; + info.numEntries = 4; + info.payloadEntries[0] = {"lastOp", ROCTX_PAYLOAD_ENTRY_TYPE_REDOP, {}}; + info.payloadEntries[0].payload.typeRedOp = static_cast(ncclNumDevRedOps - 1); + info.payloadEntries[1] = {"badOp", ROCTX_PAYLOAD_ENTRY_TYPE_REDOP, {}}; + info.payloadEntries[1].payload.typeRedOp = static_cast(ncclNumDevRedOps); + info.payloadEntries[2] = {"lastType", ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE, {}}; + info.payloadEntries[2].payload.typeDataType = static_cast(ncclNumTypes - 1); + info.payloadEntries[3] = {"badType", ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE, {}}; + info.payloadEntries[3].payload.typeDataType = static_cast(ncclNumTypes); + stringify(&info); + EXPECT_STREQ(info.message, "{X: lastOp: SumPostDiv, badOp: unknown, lastType: b8, badType: unknown}"); + roctxFree(&info); +} + +TEST_F(RoctxMicrotest, Stringify_UnmappedRoctxTypeHitsTheDefaultBranch) { + roctxPayloadInfo info; + roctxAlloc(&info, 1); + info.id = "X"; + info.numEntries = 1; + info.payloadEntries[0] = {"weird", static_cast(ROCTX_PAYLOAD_NUM_ENTRY_TYPES), {}}; + stringify(&info); + EXPECT_STREQ(info.message, "{X: weird: unknown roctx payload type}"); + roctxFree(&info); +} + +TEST_F(RoctxMicrotest, Stringify_MissingBufferIsANoOp) { + roctxPayloadInfo info; // payloadEntries and message both null + info.id = "X"; + info.numEntries = 0; + stringify(&info); // must not dereference the null message + EXPECT_EQ(info.message, nullptr); +} + +// Entries present but no message buffer: still a no-op, checked separately from the +// both-null case above because the guard is `!entries || !message`. +TEST_F(RoctxMicrotest, Stringify_EntriesWithoutMessageIsANoOp) { + roctxPayloadSchemaEntryInfo entry{"n", ROCTX_PAYLOAD_ENTRY_TYPE_INT, {}}; + roctxPayloadInfo info; + info.payloadEntries = &entry; + info.id = "X"; + info.numEntries = 1; + stringify(&info); + EXPECT_EQ(info.message, nullptr); +} + +// A mapping to a roctx type outside the enum reaches extractPayloadInfo's default branch, +// which stores the type but no payload, and stringify then prints the unknown-type text. +TEST_F(RoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefaultBranch) { + const uint64_t nvtxType = NVTX_PAYLOAD_ENTRY_TYPE_INT + 2; + ASSERT_EQ(nvtxToRoctx.count(nvtxType), 0u); + const auto bogus = static_cast(ROCTX_PAYLOAD_NUM_ENTRY_TYPES); + nvtxToRoctx[nvtxType] = bogus; + int value = 5; + nvtxPayloadSchemaEntry_t schema[] = {Entry(nvtxType, "w", 0)}; + nvtxPayloadData_t data{0, sizeof(value), &value}; + roctxPayloadInfo info; + roctxAlloc(&info, 1); + extractPayloadInfo(schema, &data, 1, "D", &info); + nvtxToRoctx.erase(nvtxType); + EXPECT_EQ(info.payloadEntries[0].type, bogus); + EXPECT_STREQ(info.message, "{D: w: unknown roctx payload type}"); + roctxFree(&info); +} + +// Pinned as observed: an nvtx type with no mapping is looked up with std::map::operator[], +// which inserts it as ROCTX_PAYLOAD_ENTRY_TYPE_INT (value 0) and decodes the field as an int. +TEST_F(RoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInsertedIntoTheMap) { + const uint64_t unmapped = NVTX_PAYLOAD_ENTRY_TYPE_INT + 1; + ASSERT_EQ(nvtxToRoctx.count(unmapped), 0u); + int value = 42; + nvtxPayloadSchemaEntry_t schema[] = {Entry(unmapped, "v", 0)}; + nvtxPayloadData_t data{0, sizeof(value), &value}; + roctxPayloadInfo info; + roctxAlloc(&info, 1); + extractPayloadInfo(schema, &data, 1, "U", &info); + EXPECT_EQ(info.payloadEntries[0].type, ROCTX_PAYLOAD_ENTRY_TYPE_INT); + EXPECT_STREQ(info.message, "{U: v: 42}"); + EXPECT_EQ(nvtxToRoctx.count(unmapped), 1u); + nvtxToRoctx.erase(unmapped); // keep the global table as other tests expect it + roctxFree(&info); +} + +// --------------------------------------------------------------------------- +// roctx_scoped_range_in +// --------------------------------------------------------------------------- + +TEST_F(RoctxMicrotest, LogRoctxParam_IsReadFromRcclLogRoctxWithDefaultZero) { + std::string seenEnv; + int64_t seenDefault = -1; + g_loadParam = [&](const char* env, int64_t deftVal) { + seenEnv = env; + seenDefault = deftVal; + return deftVal; + }; + EXPECT_EQ(rcclParamLogRoctx(), 0); + EXPECT_EQ(seenEnv, "RCCL_LOG_ROCTX"); + EXPECT_EQ(seenDefault, 0); +} + +TEST_F(RoctxMicrotest, ScopedRange_LoggingOff_DoesNothing) { + SetLogRoctx(0); + { + roctx_scoped_range_in range("msg"); + roctx_scoped_range_in defaulted; + } + EXPECT_TRUE(g_logged.empty()); +} + +// Without ROCTX_ENABLE each push and each pop falls back to a WARN, so one range is two WARNs. +TEST_F(RoctxMicrotest, ScopedRange_LoggingOnWithoutRoctx_WarnsOnEnterAndExit) { + SetLogRoctx(1); + { + roctx_scoped_range_in range("msg"); + EXPECT_EQ(WarnCount(), 1u); + } + EXPECT_EQ(WarnCount(), 2u); + for (const auto& m : g_logged) EXPECT_NE(m.text.find("ROCTX_ENABLE is not defined"), std::string::npos); +} + +TEST_F(RoctxMicrotest, ScopedRange_DefaultConstructorDelegatesToTheMessageConstructor) { + SetLogRoctx(1); + { roctx_scoped_range_in range; } + EXPECT_EQ(WarnCount(), 2u); +} + +// The payload constructor builds (and the destructor frees) the message even though it is +// never pushed on this build; LeakSanitizer would flag a missing roctxFree. +TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOn_WarnsOnEnterAndExit) { + SetLogRoctx(1); + MixedPayload payload{1, 2, ncclDevSum, ncclInt8}; + nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "n", offsetof(MixedPayload, i))}; + nvtxPayloadData_t data{0, sizeof(payload), &payload}; + { + roctx_scoped_range_in range(schema, &data, 1, "RCCL_AllReduce"); + EXPECT_EQ(WarnCount(), 1u); + } + EXPECT_EQ(WarnCount(), 2u); +} + +TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOff_DoesNothing) { + SetLogRoctx(0); + int value = 1; + nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "n", 0)}; + nvtxPayloadData_t data{0, sizeof(value), &value}; + { roctx_scoped_range_in range(schema, &data, 1, "RCCL_AllReduce"); } + EXPECT_TRUE(g_logged.empty()); +} + +// --------------------------------------------------------------------------- +// nvmlwrap_stub.cc: no NVML on ROCm, so every query reports "no device". +// --------------------------------------------------------------------------- + +namespace { +nvmlDevice_t FakeDevice() { return reinterpret_cast(0x1234); } +} // namespace + +TEST(NvmlStubMicrotest, LifecycleCallsSucceed) { + EXPECT_EQ(ncclNvmlSymbols(), ncclSuccess); + EXPECT_EQ(ncclNvmlInit(), ncclSuccess); + EXPECT_EQ(ncclNvmlShutdown(), ncclSuccess); +} + +TEST(NvmlStubMicrotest, ReportsZeroDevices) { + EXPECT_EQ(ncclNvmlDeviceCount, 0); + unsigned int count = 99; + EXPECT_EQ(ncclNvmlDeviceGetCount(&count), ncclSuccess); + EXPECT_EQ(count, 0u); +} + +// The successful getters write a fixed value; check it lands rather than a stale input. +TEST(NvmlStubMicrotest, SucceedingGettersWriteFixedValues) { + unsigned index = 99; + EXPECT_EQ(ncclNvmlDeviceGetIndex(FakeDevice(), &index), ncclSuccess); + EXPECT_EQ(index, 0u); + + unsigned int minor = 99; + EXPECT_EQ(ncclNvmlDeviceGetMinorNumber(FakeDevice(), &minor), ncclSuccess); + EXPECT_EQ(minor, 0u); + + int major = 99, minorCc = 99; + EXPECT_EQ(ncclNvmlDeviceGetCudaComputeCapability(FakeDevice(), &major, &minorCc), ncclSuccess); + EXPECT_EQ(major, 1); + EXPECT_EQ(minorCc, 1); +} + +// The failing lookups return ncclSystemError and leave their out-params alone. +TEST(NvmlStubMicrotest, DeviceLookupsFailWithoutWritingOutputs) { + nvmlDevice_t sentinel = FakeDevice(); + nvmlDevice_t dev = sentinel; + EXPECT_EQ(ncclNvmlDeviceGetHandleByPciBusId("0000:00:00.0", &dev), ncclSystemError); + EXPECT_EQ(dev, sentinel); + EXPECT_EQ(ncclNvmlDeviceGetHandleByIndex(0, &dev), ncclSystemError); + EXPECT_EQ(dev, sentinel); + + nvmlPciInfo_t pci; + std::memset(&pci, 0xAB, sizeof(pci)); + EXPECT_EQ(ncclNvmlDeviceGetPciInfo(FakeDevice(), &pci), ncclSystemError); + EXPECT_EQ(ncclNvmlDeviceGetNvLinkRemotePciInfo(FakeDevice(), 0, &pci), ncclSystemError); + const unsigned char* bytes = reinterpret_cast(&pci); + for (size_t i = 0; i < sizeof(pci); ++i) ASSERT_EQ(bytes[i], 0xAB) << "byte " << i; + + nvmlEnableState_t state = NVML_FEATURE_ENABLED; + EXPECT_EQ(ncclNvmlDeviceGetNvLinkState(FakeDevice(), 0, &state), ncclSystemError); + EXPECT_EQ(state, NVML_FEATURE_ENABLED); + + unsigned int value = 99; + EXPECT_EQ(ncclNvmlDeviceGetNvLinkCapability(FakeDevice(), 0, NVML_NVLINK_CAP_P2P_SUPPORTED, &value), + ncclSystemError); + EXPECT_EQ(ncclNvmlDeviceGetCurrPcieLinkGeneration(FakeDevice(), &value), ncclSystemError); + EXPECT_EQ(ncclNvmlDeviceGetCurrPcieLinkWidth(FakeDevice(), &value), ncclSystemError); + EXPECT_EQ(value, 99u); +} + +// These fail too, but clear their output first so a caller that ignores the error +// reads a defined value instead of stack garbage. +TEST(NvmlStubMicrotest, FailingQueriesClearTheirOutputs) { + char name[8] = "garbage"; + EXPECT_EQ(ncclNvmlDeviceGetName(FakeDevice(), name, sizeof(name)), ncclSystemError); + EXPECT_EQ(name[0], '\0'); + + char version[8] = "garbage"; + EXPECT_EQ(ncclNvmlSystemGetDriverVersion(version, sizeof(version)), ncclSystemError); + EXPECT_EQ(version[0], '\0'); + + unsigned long long count = 99; + EXPECT_EQ(ncclNvmlDeviceGetMemoryErrorCounter(FakeDevice(), NVML_MEMORY_ERROR_TYPE_CORRECTED, NVML_VOLATILE_ECC, + NVML_MEMORY_LOCATION_DRAM, &count), + ncclSystemError); + EXPECT_EQ(count, 0u); +} + +// A zero length means "no buffer": nothing may be written, not even the terminator. +TEST(NvmlStubMicrotest, ZeroLengthBuffersAreNotWritten) { + char name[2] = {'n', 'x'}; + EXPECT_EQ(ncclNvmlDeviceGetName(FakeDevice(), name, 0), ncclSystemError); + EXPECT_EQ(name[0], 'n'); + + char version[2] = {'v', 'x'}; + EXPECT_EQ(ncclNvmlSystemGetDriverVersion(version, 0), ncclSystemError); + EXPECT_EQ(version[0], 'v'); +} + +TEST(NvmlStubMicrotest, FieldValuesMarksEveryRequestedFieldNotSupported) { + nvmlFieldValue_t values[3]; + std::memset(values, 0, sizeof(values)); + for (auto& v : values) v.nvmlReturn = NVML_SUCCESS; + EXPECT_EQ(ncclNvmlDeviceGetFieldValues(FakeDevice(), 2, values), ncclSystemError); + EXPECT_EQ(values[0].nvmlReturn, NVML_ERROR_NOT_SUPPORTED); + EXPECT_EQ(values[1].nvmlReturn, NVML_ERROR_NOT_SUPPORTED); + EXPECT_EQ(values[2].nvmlReturn, NVML_SUCCESS) << "only valuesCount entries may be touched"; + + EXPECT_EQ(ncclNvmlDeviceGetFieldValues(FakeDevice(), 0, values), ncclSystemError); +} diff --git a/projects/rccl/test/host/run_host_tests.sh b/projects/rccl/test/host/run_host_tests.sh index ec702f4f62b1..8047a8c653d6 100755 --- a/projects/rccl/test/host/run_host_tests.sh +++ b/projects/rccl/test/host/run_host_tests.sh @@ -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-UnitTestsMicroRoctxNvml:$SCRIPT_DIR/host_tests_micro_roctx_nvml.xml" "rccl-UnitTestsMicroGinHost:$SCRIPT_DIR/host_tests_micro_gin_host.xml" "rccl-UnitTestsMicroDiagnostics:$SCRIPT_DIR/host_tests_micro_diagnostics.xml" ) diff --git a/projects/rccl/test/test_categories_micro.yaml b/projects/rccl/test/test_categories_micro.yaml index a940ce5f4f7c..566e6d1f387a 100644 --- a/projects/rccl/test/test_categories_micro.yaml +++ b/projects/rccl/test/test_categories_micro.yaml @@ -114,6 +114,9 @@ test_categories: # ce_coll.cc (rccl-UnitTestsMicroCe) - "CeCollMicrotest.*" - "LaunchPaths/CeAlltoAllvSyncMicrotest.*" + # misc/roctx.cc, misc/nvmlwrap_stub.cc (rccl-UnitTestsMicroRoctxNvml) + - "RoctxMicrotest.*" + - "NvmlStubMicrotest.*" # dev_runtime.cc - "Alloc*" - "Comm*" From c15b079725875ded3d9edca5ec1d41d9eb602d83 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Tue, 13 Oct 2026 09:40:38 -0700 Subject: [PATCH 04/16] test(rccl): cover roctx.cc's ROCTX_ENABLE arm in the roctx microtests In-tree builds copy the rccl target's COMPILE_DEFINITIONS (PRIVATE ones included) into every test target, so ROCTX_ENABLE reaches this binary whenever roctracer is found. roctx.cc then calls roctxRangePushA/Pop, which nothing on the micro-test link line provides; that was the precheckin "Failed in Compile stage". The standalone build never defines it, so this did not reproduce locally. Define both entry points in the test TU to record calls, and make the scoped-range tests hold in either build: push/pop events and the exact message pushed with ROCTX_ENABLE, the WARN fallback without it. Co-Authored-By: Claude --- .../rccl/test/host/roctx-nvml-stub-test.cc | 91 ++++++++++++++++--- 1 file changed, 76 insertions(+), 15 deletions(-) diff --git a/projects/rccl/test/host/roctx-nvml-stub-test.cc b/projects/rccl/test/host/roctx-nvml-stub-test.cc index 6f375cec7e91..4c41281ca501 100644 --- a/projects/rccl/test/host/roctx-nvml-stub-test.cc +++ b/projects/rccl/test/host/roctx-nvml-stub-test.cc @@ -14,9 +14,11 @@ // g_loadParam (param_redirect.h routes rcclParamLogRoctx() through it on every // call instead of caching, so each test picks LOG_ROCTX independently). // -// The host build never defines ROCTX_ENABLE, so roctxRangePushA/Pop compile out -// and the scoped range's observable behaviour on this build is its WARN -// fallback plus the payload alloc/free it performs around it. +// The scoped range has two production arms. In-tree builds where roctracer is +// found inherit ROCTX_ENABLE from the rccl target and call roctxRangePushA/Pop, +// which this TU defines to record the messages pushed. The standalone build +// has no ROCTX_ENABLE and falls back to one WARN per enter/exit. The range +// tests are written to hold in either build. // // Every stringify() case here stays under MAX_MESSAGE_LENGTH on purpose: past // it, stringify() writes beyond its 1024-byte buffer (confirmed under ASan), @@ -64,6 +66,26 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long /*flags*/, const char* namespace { std::function g_roctxMalloc; } // namespace + +// In-tree builds inherit the rccl target's COMPILE_DEFINITIONS (test/CMakeLists.txt), so +// ROCTX_ENABLE is defined here whenever roctracer was found, and roctx.cc then calls the real +// roctracer entry points. Nothing on the micro-test link line provides them, so record the calls +// instead; that also lets the range tests check the exact message pushed. The standalone build has +// no ROCTX_ENABLE and exercises roctx.cc's WARN fallback instead. +#ifdef ROCTX_ENABLE +namespace { +std::vector g_roctxPushed; +int g_roctxPops = 0; +} // namespace +int roctxRangePushA(const char* message) { + g_roctxPushed.push_back(message ? message : ""); + return 0; +} +int roctxRangePop() { + ++g_roctxPops; + return 0; +} +#endif static void* RoctxTestMalloc(size_t n) { return g_roctxMalloc ? g_roctxMalloc(n) : std::malloc(n); } #define malloc(n) RoctxTestMalloc(n) #include ROCTX_CC_PATH @@ -75,11 +97,40 @@ namespace { class RoctxMicrotest : public ::testing::Test { protected: - void SetUp() override { g_logged.clear(); } + void SetUp() override { ResetRecorders(); } void TearDown() override { g_loadParam = [](const char*, int64_t deftVal) { return deftVal; }; g_roctxMalloc = nullptr; + ResetRecorders(); + } + + static void ResetRecorders() { g_logged.clear(); +#ifdef ROCTX_ENABLE + g_roctxPushed.clear(); + g_roctxPops = 0; +#endif + } + + // Range enter/exit events, whichever form this build gives them: roctracer push/pop calls with + // ROCTX_ENABLE, one WARN each without it. With logging off there must be none either way. + static size_t RangeEvents() { +#ifdef ROCTX_ENABLE + EXPECT_TRUE(g_logged.empty()) << "a ROCTX_ENABLE build must not fall back to WARN"; + return g_roctxPushed.size() + static_cast(g_roctxPops); +#else + return WarnCount(); +#endif + } + + // ROCTX_ENABLE builds can also check the message handed to roctracer; the fallback never sees it. + static void ExpectPushedMessage(const char* expected) { +#ifdef ROCTX_ENABLE + ASSERT_EQ(g_roctxPushed.size(), 1u); + EXPECT_EQ(g_roctxPushed[0], expected); +#else + (void)expected; +#endif } // Lets the first `succeedFirst` roctx.cc mallocs through, then fails every one after. @@ -378,38 +429,47 @@ TEST_F(RoctxMicrotest, ScopedRange_LoggingOff_DoesNothing) { roctx_scoped_range_in range("msg"); roctx_scoped_range_in defaulted; } + EXPECT_EQ(RangeEvents(), 0u); EXPECT_TRUE(g_logged.empty()); } -// Without ROCTX_ENABLE each push and each pop falls back to a WARN, so one range is two WARNs. -TEST_F(RoctxMicrotest, ScopedRange_LoggingOnWithoutRoctx_WarnsOnEnterAndExit) { +// One range is one event on enter and one on exit: a roctracer push/pop with ROCTX_ENABLE, +// otherwise the WARN that stands in for each. +TEST_F(RoctxMicrotest, ScopedRange_LoggingOn_OneEventOnEnterAndOneOnExit) { SetLogRoctx(1); { roctx_scoped_range_in range("msg"); - EXPECT_EQ(WarnCount(), 1u); + EXPECT_EQ(RangeEvents(), 1u); + ExpectPushedMessage("msg"); } - EXPECT_EQ(WarnCount(), 2u); + EXPECT_EQ(RangeEvents(), 2u); +#ifndef ROCTX_ENABLE for (const auto& m : g_logged) EXPECT_NE(m.text.find("ROCTX_ENABLE is not defined"), std::string::npos); +#endif } TEST_F(RoctxMicrotest, ScopedRange_DefaultConstructorDelegatesToTheMessageConstructor) { SetLogRoctx(1); - { roctx_scoped_range_in range; } - EXPECT_EQ(WarnCount(), 2u); + { + roctx_scoped_range_in range; + ExpectPushedMessage(""); + } + EXPECT_EQ(RangeEvents(), 2u); } -// The payload constructor builds (and the destructor frees) the message even though it is -// never pushed on this build; LeakSanitizer would flag a missing roctxFree. -TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOn_WarnsOnEnterAndExit) { +// The payload constructor builds the message from the schema (pushed verbatim with ROCTX_ENABLE) +// and the destructor frees it in both builds; LeakSanitizer would flag a missing roctxFree. +TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOn_PushesTheStringifiedPayload) { SetLogRoctx(1); MixedPayload payload{1, 2, ncclDevSum, ncclInt8}; nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "n", offsetof(MixedPayload, i))}; nvtxPayloadData_t data{0, sizeof(payload), &payload}; { roctx_scoped_range_in range(schema, &data, 1, "RCCL_AllReduce"); - EXPECT_EQ(WarnCount(), 1u); + EXPECT_EQ(RangeEvents(), 1u); + ExpectPushedMessage("{RCCL_AllReduce: n: 1}"); } - EXPECT_EQ(WarnCount(), 2u); + EXPECT_EQ(RangeEvents(), 2u); } TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOff_DoesNothing) { @@ -418,6 +478,7 @@ TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOff_DoesNothing) nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "n", 0)}; nvtxPayloadData_t data{0, sizeof(value), &value}; { roctx_scoped_range_in range(schema, &data, 1, "RCCL_AllReduce"); } + EXPECT_EQ(RangeEvents(), 0u); EXPECT_TRUE(g_logged.empty()); } From c9c918a1ca6c0c7a5ec8e44ab6a03c5618cc65be Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Tue, 13 Oct 2026 12:38:05 -0700 Subject: [PATCH 05/16] test(rccl): build roctx.cc's ROCTX_ENABLE arm standalone, split push/pop checks, address ARGUS round 2 Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 17 +- projects/rccl/test/host/MICROTEST_README.md | 25 +- .../fakes/roctx_include/roctracer/roctx.h | 27 ++ .../rccl/test/host/roctx-nvml-stub-test.cc | 236 +++++++++++------- projects/rccl/test/host/run_host_tests.sh | 3 + projects/rccl/test/test_categories_micro.yaml | 5 +- 6 files changed, 208 insertions(+), 105 deletions(-) create mode 100644 projects/rccl/test/host/fakes/roctx_include/roctracer/roctx.h diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 9621438ef4cf..bd779d79e25b 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -753,6 +753,9 @@ function(rccl_standalone_setup) # binary; run_host_tests.sh (the only in-repo workflow that runs these) uses # this standalone path, so build the shipping arm explicitly too. set(_ENQ_BUILD_DEVLINKER TRUE PARENT_SCOPE) + # Same reason for roctx.cc's ROCTX_ENABLE arm: in-tree inherits it from the rccl target when + # roctracer is found, so standalone builds that arm explicitly. + set(_MICRO_BUILD_ROCTX_VARIANT TRUE PARENT_SCOPE) # No GENERATE_SYM_KERNELS variable exists in this standalone project(); check # for the generated host TU on disk instead. if(EXISTS "${RCCL_HIPIFY_DIR}/gensrc/symmetric/sym_kernels_host.cc") @@ -789,6 +792,7 @@ if(BUILD_TESTS) set(_MICRO_MODE_DEFS "") set(_MICRO_AMDSMI_DEF "") set(_ENQ_BUILD_DEVLINKER FALSE) + set(_MICRO_BUILD_ROCTX_VARIANT FALSE) set(_MICRO_BUILD_SYM_KERNELS ${GENERATE_SYM_KERNELS}) else() # --- hipcc is mandatory (must be resolved before project() at top level) --- @@ -1068,7 +1072,7 @@ rccl_add_micro_binary(rccl-UnitTestsMicroClientSupport # --------------------------------------------------------------------------- # rccl-UnitTestsMicroRoctxNvml # --------------------------------------------------------------------------- -rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml +set(_roctx_nvml_common_args SOURCES ${TEST_MICRO_ROCTX_NVML_SOURCE_FILES} EXTRA_DEFS ${_MICRO_MODE_DEFS} @@ -1076,6 +1080,17 @@ rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml NVMLWRAP_STUB_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/nvmlwrap_stub.cc" CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml") +rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml ${_roctx_nvml_common_args}) +# roctx.cc gates its roctracer calls on ROCTX_ENABLE at the preprocessor, so one compile covers +# one arm. The -roctx variant covers the roctracer arm against a stand-in +# (fakes/roctx_include), so it builds whether or not roctracer is installed. +if(_MICRO_BUILD_ROCTX_VARIANT) + rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml-roctx ${_roctx_nvml_common_args}) + target_compile_definitions(rccl-UnitTestsMicroRoctxNvml-roctx PRIVATE ROCTX_ENABLE) + target_include_directories(rccl-UnitTestsMicroRoctxNvml-roctx BEFORE PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/fakes/roctx_include) +endif() + # --------------------------------------------------------------------------- # rccl-UnitTestsMicroSymKernels # --------------------------------------------------------------------------- diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index 1ff07e6588f5..92709661e2bc 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -248,6 +248,19 @@ 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-UnitTestsMicroRoctxNvml`** (+ **`-roctx`**) — `src/misc/roctx.cc` and + `src/misc/nvmlwrap_stub.cc` (via `ROCTX_CC_PATH`/`NVMLWRAP_STUB_CC_PATH`, both + from `roctx-nvml-stub-test.cc`); suites `RoctxNvmlRoctxMicrotest.*` and + `RoctxNvmlStubMicrotest.*`, both caught by the `RoctxNvml*` pattern. Its own + binary because `fakes/comm_fakes.cc`, linked by most other targets, stubs the + `roctx_scoped_range_in` constructor/destructor this file compiles for real, + and `ras-diagnostics-gpu-test.cc` defines its own NVML seams; the TU supplies + its own `ncclDebugLog` and `g_loadParam`, so no shared fakes are linked. + `roctx.cc` gates its `roctxRangePushA`/`roctxRangePop` calls on + `ROCTX_ENABLE` at the preprocessor: the base binary covers the WARN fallback, + and the standalone-only `-roctx` variant defines `ROCTX_ENABLE` against the + stand-in `fakes/roctx_include/roctracer/roctx.h`, with the test TU recording + each push and pop. - **`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.*`, @@ -829,19 +842,21 @@ 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-UnitTestsMicroRoctxNvml[-roctx]`, `rccl-UnitTestsMicroSymKernels` and `rccl-UnitTestsMicroTaskPrep` — **without configuring/building all of librccl**. It compiles just the tests + fakes + the hipified unit-under-test sources. -Three of those names are preprocessor variants, not duplicates. `init.cc` gates +Four of those names are preprocessor variants, not duplicates. `init.cc` gates part of its allocation path on `HIP_*_UNCACHED_MEMORY` and its fault-mask blocks on `ENABLE_FAULT_INJECTION`, and `enqueue.cc` gates `rcclShmemDynamicSize` on -`RCCL_DEVICE_LINKER`, all at the **preprocessor**, so one compile can only ever -reach one arm. The in-RCCL-build path inherits `RCCL_DEVICE_LINKER` from the +`RCCL_DEVICE_LINKER`, and `roctx.cc` gates its roctracer calls on `ROCTX_ENABLE`, +all at the **preprocessor**, so one compile can only ever reach one arm. The in-RCCL-build path inherits `RCCL_DEVICE_LINKER` from the `rccl` target's compile definitions (`ENABLE_DEVICE_LINKER` defaults ON, so the device-linker arm is the one that ships); this standalone project has no `rccl` target to inherit from, which is why it builds the `-devlinker` variant -explicitly. `ENABLE_FAULT_INJECTION` was the same inheritance in reverse - +explicitly, and likewise the `-roctx` variant for `ROCTX_ENABLE`, which in-tree +is inherited only when roctracer is found. `ENABLE_FAULT_INJECTION` was the same inheritance in reverse - defined in-tree, absent standalone, so the two paths tested different code - and is now stripped from the inherited list and set per variant in both paths, so the `-faultinj` binary is the only one that has it either way. @@ -868,6 +883,8 @@ cmake --build build -j"$(nproc)" ./build/rccl-UnitTestsMicroSymKernels # sym_kernels.cc tests ./build/rccl-UnitTestsMicroTaskPrep # src/enqueue/task_prep/ + task_sched/ tests ./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/{p2p,ib_write_bw}.cc tests +./build/rccl-UnitTestsMicroRoctxNvml # misc/roctx.cc + misc/nvmlwrap_stub.cc tests +./build/rccl-UnitTestsMicroRoctxNvml-roctx # same, ROCTX_ENABLE arm ./build/rccl-UnitTestsMicroGinHost # src/gin/gin_host.cc GIN_PROXY_NTHREADS ./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/p2p.cc tests ./build/rccl-UnitTestsMicroCe # src/ce_coll.cc tests diff --git a/projects/rccl/test/host/fakes/roctx_include/roctracer/roctx.h b/projects/rccl/test/host/fakes/roctx_include/roctracer/roctx.h new file mode 100644 index 000000000000..553f30c5448a --- /dev/null +++ b/projects/rccl/test/host/fakes/roctx_include/roctracer/roctx.h @@ -0,0 +1,27 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Stand-in for roctracer's , used only by the standalone +// rccl-UnitTestsMicroRoctxNvml-roctx target so roctx.cc's ROCTX_ENABLE arm +// compiles whether or not roctracer is installed. Same C-linkage declarations +// as the real header for the two entry points roctx.cc calls; +// roctx-nvml-stub-test.cc defines them to record each call. + +#ifndef RCCL_TEST_HOST_FAKES_ROCTRACER_ROCTX_H_ +#define RCCL_TEST_HOST_FAKES_ROCTRACER_ROCTX_H_ + +#ifdef __cplusplus +extern "C" { +#endif + +int roctxRangePushA(const char* message); +int roctxRangePop(void); + +#ifdef __cplusplus +} +#endif + +#endif // RCCL_TEST_HOST_FAKES_ROCTRACER_ROCTX_H_ diff --git a/projects/rccl/test/host/roctx-nvml-stub-test.cc b/projects/rccl/test/host/roctx-nvml-stub-test.cc index 4c41281ca501..c5a295c72315 100644 --- a/projects/rccl/test/host/roctx-nvml-stub-test.cc +++ b/projects/rccl/test/host/roctx-nvml-stub-test.cc @@ -14,11 +14,13 @@ // g_loadParam (param_redirect.h routes rcclParamLogRoctx() through it on every // call instead of caching, so each test picks LOG_ROCTX independently). // -// The scoped range has two production arms. In-tree builds where roctracer is -// found inherit ROCTX_ENABLE from the rccl target and call roctxRangePushA/Pop, -// which this TU defines to record the messages pushed. The standalone build -// has no ROCTX_ENABLE and falls back to one WARN per enter/exit. The range -// tests are written to hold in either build. +// The scoped range has two production arms. With ROCTX_ENABLE it calls +// roctxRangePushA/Pop, which this TU defines to record the messages pushed; +// without it, it falls back to one WARN per enter/exit. In-tree builds inherit +// ROCTX_ENABLE from the rccl target when roctracer is found. The standalone +// build compiles this TU twice: rccl-UnitTestsMicroRoctxNvml without it, and +// rccl-UnitTestsMicroRoctxNvml-roctx with it against the stand-in +// fakes/roctracer/roctx.h, so CI runs both arms. The range tests hold in either. // // Every stringify() case here stays under MAX_MESSAGE_LENGTH on purpose: past // it, stringify() writes beyond its 1024-byte buffer (confirmed under ASan), @@ -31,6 +33,7 @@ #include #include #include +#include #include #include @@ -67,11 +70,9 @@ namespace { std::function g_roctxMalloc; } // namespace -// In-tree builds inherit the rccl target's COMPILE_DEFINITIONS (test/CMakeLists.txt), so -// ROCTX_ENABLE is defined here whenever roctracer was found, and roctx.cc then calls the real -// roctracer entry points. Nothing on the micro-test link line provides them, so record the calls -// instead; that also lets the range tests check the exact message pushed. The standalone build has -// no ROCTX_ENABLE and exercises roctx.cc's WARN fallback instead. +// With ROCTX_ENABLE (inherited in-tree when roctracer is found, or set by the standalone -roctx +// variant), roctx.cc calls the roctracer entry points. Nothing on the micro-test link line provides +// them, so record the calls instead; that also lets the range tests check the exact message pushed. #ifdef ROCTX_ENABLE namespace { std::vector g_roctxPushed; @@ -95,12 +96,18 @@ static void* RoctxTestMalloc(size_t n) { return g_roctxMalloc ? g_roctxMalloc(n) namespace { -class RoctxMicrotest : public ::testing::Test { +class RoctxNvmlRoctxMicrotest : public ::testing::Test { protected: - void SetUp() override { ResetRecorders(); } + void SetUp() override { + ResetRecorders(); + savedNvtxToRoctx_ = nvtxToRoctx; + } void TearDown() override { g_loadParam = [](const char*, int64_t deftVal) { return deftVal; }; g_roctxMalloc = nullptr; + // nvtxToRoctx is a production global that two tests insert into; restoring it here rather + // than inline keeps a failed assertion in one of them from leaking into the next test. + nvtxToRoctx = savedNvtxToRoctx_; ResetRecorders(); } @@ -133,6 +140,29 @@ class RoctxMicrotest : public ::testing::Test { #endif } + // After a range's scope ends: exactly one enter and one exit, checked separately so a + // destructor that pushed instead of popping cannot pass as two events. + static void ExpectOneRangeEnteredAndExited() { +#ifdef ROCTX_ENABLE + EXPECT_TRUE(g_logged.empty()) << "a ROCTX_ENABLE build must not fall back to WARN"; + EXPECT_EQ(g_roctxPushed.size(), 1u); + EXPECT_EQ(g_roctxPops, 1); +#else + EXPECT_EQ(WarnCount(), 2u); +#endif + } + + // Counts roctx.cc's mallocs while letting every one through. + static int* CountMallocs() { + static int calls; + calls = 0; + g_roctxMalloc = [](size_t n) -> void* { + ++calls; + return std::malloc(n); + }; + return &calls; + } + // Lets the first `succeedFirst` roctx.cc mallocs through, then fails every one after. static int* FailMallocAfter(int succeedFirst) { static int calls; @@ -153,6 +183,25 @@ class RoctxMicrotest : public ::testing::Test { return n; } + // Runs stringify() over `entries` under `id` and returns the message it built. + static std::string Stringify(const char* id, const std::vector& entries) { + roctxPayloadInfo info; + roctxAlloc(&info, entries.empty() ? 1 : entries.size()); + info.id = id; + info.numEntries = entries.size(); + for (size_t i = 0; i < entries.size(); ++i) info.payloadEntries[i] = entries[i]; + stringify(&info); + std::string message = info.message; + roctxFree(&info); + return message; + } + + static roctxPayloadSchemaEntryInfo IntEntry(const char* name, int value) { + roctxPayloadSchemaEntryInfo e{name, ROCTX_PAYLOAD_ENTRY_TYPE_INT, {}}; + e.payload.typeInt = value; + return e; + } + // Zero-initialised schema entry: the fields roctx.cc does not read stay at 0. static nvtxPayloadSchemaEntry_t Entry(uint64_t type, const char* name, uint64_t offset) { nvtxPayloadSchemaEntry_t e; @@ -162,6 +211,9 @@ class RoctxMicrotest : public ::testing::Test { e.offset = offset; return e; } + +private: + std::map savedNvtxToRoctx_; }; struct MixedPayload { @@ -177,7 +229,7 @@ struct MixedPayload { // Lookup tables // --------------------------------------------------------------------------- -TEST_F(RoctxMicrotest, NvtxToRoctx_MapsEachSupportedNvtxTypeToItsRoctxType) { +TEST_F(RoctxNvmlRoctxMicrotest, NvtxToRoctx_MapsEachSupportedNvtxTypeToItsRoctxType) { EXPECT_EQ(nvtxToRoctx.size(), 4u); EXPECT_EQ(nvtxToRoctx.at(NVTX_PAYLOAD_ENTRY_TYPE_INT), ROCTX_PAYLOAD_ENTRY_TYPE_INT); EXPECT_EQ(nvtxToRoctx.at(NVTX_PAYLOAD_ENTRY_TYPE_SIZE), ROCTX_PAYLOAD_ENTRY_TYPE_SIZE); @@ -186,7 +238,7 @@ TEST_F(RoctxMicrotest, NvtxToRoctx_MapsEachSupportedNvtxTypeToItsRoctxType) { } // The names stringify() prints for redop/datatype payloads, indexed by enum value. -TEST_F(RoctxMicrotest, NameTables_MatchTheEnumOrder) { +TEST_F(RoctxNvmlRoctxMicrotest, NameTables_MatchTheEnumOrder) { const char* redOps[] = {"Sum", "Prod", "MinMax", "PreMulSum", "SumPostDiv"}; ASSERT_EQ(sizeof(redOps) / sizeof(redOps[0]), static_cast(ncclNumDevRedOps)); for (int i = 0; i < ncclNumDevRedOps; ++i) EXPECT_STREQ(ncclRedOpStr[i], redOps[i]) << "op " << i; @@ -196,32 +248,33 @@ TEST_F(RoctxMicrotest, NameTables_MatchTheEnumOrder) { for (int i = 0; i < ncclNumTypes; ++i) EXPECT_STREQ(ncclDataTypeStr[i], dataTypes[i]) << "type " << i; } -// Pinned as observed, not endorsed: the table has 4 slots but only 3 initialisers, so the -// DATATYPE slot is null. Nothing in the tree reads roctxEntryTypeStr today. -TEST_F(RoctxMicrotest, RoctxEntryTypeStr_DatatypeSlotIsNull) { +// One slot per roctx entry type. The DATATYPE slot has no initialiser today, but nothing in the +// tree reads this table, so that gap is left unpinned rather than locked in. +TEST_F(RoctxNvmlRoctxMicrotest, RoctxEntryTypeStr_HasOneSlotPerEntryType) { + static_assert(sizeof(roctxEntryTypeStr) / sizeof(roctxEntryTypeStr[0]) == ROCTX_PAYLOAD_NUM_ENTRY_TYPES, + "roctxEntryTypeStr must have one slot per roctxPayloadEntryType"); EXPECT_STREQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_INT], "ROCTX_PAYLOAD_ENTRY_TYPE_INT"); EXPECT_STREQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_SIZE], "ROCTX_PAYLOAD_ENTRY_TYPE_SIZE"); EXPECT_STREQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_REDOP], "ROCTX_PAYLOAD_ENTRY_TYPE_REDOP"); - EXPECT_EQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE], nullptr); } // --------------------------------------------------------------------------- // roctxAlloc / roctxFree // --------------------------------------------------------------------------- -TEST_F(RoctxMicrotest, Alloc_AllocatesEntriesAndMessage_FreeReleasesBoth) { +TEST_F(RoctxNvmlRoctxMicrotest, Alloc_AllocatesEntriesAndMessage_FreeReleasesBoth) { roctxPayloadInfo info; roctxAlloc(&info, 3); ASSERT_NE(info.payloadEntries, nullptr); ASSERT_NE(info.message, nullptr); - // Both buffers are usable at their full advertised size (ASan flags a short allocation). + // Both buffers are usable at their full advertised size (under ASan, a short allocation fails here). std::memset(info.payloadEntries, 0, 3 * sizeof(roctxPayloadSchemaEntryInfo)); std::memset(info.message, 'x', MAX_MESSAGE_LENGTH); roctxFree(&info); } // The entries allocation fails: both pointers end up null and the message is never allocated. -TEST_F(RoctxMicrotest, Alloc_EntriesAllocationFails_LeavesBothPointersNull) { +TEST_F(RoctxNvmlRoctxMicrotest, Alloc_EntriesAllocationFails_LeavesBothPointersNull) { int* calls = FailMallocAfter(0); roctxPayloadInfo info; info.message = reinterpret_cast(0x1); // must be overwritten with nullptr @@ -232,8 +285,9 @@ TEST_F(RoctxMicrotest, Alloc_EntriesAllocationFails_LeavesBothPointersNull) { } // The message allocation fails after the entries succeeded: the entries buffer is released -// (LeakSanitizer catches a missing free) and nulled, so a later roctxFree cannot double-free. -TEST_F(RoctxMicrotest, Alloc_MessageAllocationFails_ReleasesEntriesAndNullsBoth) { +// (under ASan, LeakSanitizer catches a missing free) and nulled, so a later roctxFree cannot +// double-free. +TEST_F(RoctxNvmlRoctxMicrotest, Alloc_MessageAllocationFails_ReleasesEntriesAndNullsBoth) { int* calls = FailMallocAfter(1); roctxPayloadInfo info; roctxAlloc(&info, 2); @@ -243,7 +297,7 @@ TEST_F(RoctxMicrotest, Alloc_MessageAllocationFails_ReleasesEntriesAndNullsBoth) roctxFree(&info); } -TEST_F(RoctxMicrotest, Free_ToleratesNullPointers) { +TEST_F(RoctxNvmlRoctxMicrotest, Free_ToleratesNullPointers) { roctxPayloadInfo info; // both members default to nullptr roctxFree(&info); EXPECT_EQ(info.payloadEntries, nullptr); @@ -254,7 +308,7 @@ TEST_F(RoctxMicrotest, Free_ToleratesNullPointers) { // extractPayloadInfo / stringify // --------------------------------------------------------------------------- -TEST_F(RoctxMicrotest, Extract_NoEntriesBuffer_ReturnsWithoutTouchingInfo) { +TEST_F(RoctxNvmlRoctxMicrotest, Extract_NoEntriesBuffer_ReturnsWithoutTouchingInfo) { roctxPayloadInfo info; info.id = "untouched"; info.numEntries = 77; @@ -266,7 +320,7 @@ TEST_F(RoctxMicrotest, Extract_NoEntriesBuffer_ReturnsWithoutTouchingInfo) { EXPECT_EQ(info.numEntries, 77u); } -TEST_F(RoctxMicrotest, Extract_ReadsEachTypeAtItsOffsetAndStringifies) { +TEST_F(RoctxNvmlRoctxMicrotest, Extract_ReadsEachTypeAtItsOffsetAndStringifies) { MixedPayload payload{-3, size_t{1} << 40, ncclDevProd, ncclFloat32}; nvtxPayloadSchemaEntry_t schema[] = { Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "count", offsetof(MixedPayload, i)), @@ -295,61 +349,35 @@ TEST_F(RoctxMicrotest, Extract_ReadsEachTypeAtItsOffsetAndStringifies) { } // The separator is only written between entries: one entry has no trailing ", ". -TEST_F(RoctxMicrotest, Stringify_SingleEntryHasNoSeparator) { - roctxPayloadInfo info; - roctxAlloc(&info, 1); - info.id = "Bcast"; - info.numEntries = 1; - info.payloadEntries[0].name = "root"; - info.payloadEntries[0].type = ROCTX_PAYLOAD_ENTRY_TYPE_INT; - info.payloadEntries[0].payload.typeInt = 7; - stringify(&info); - EXPECT_STREQ(info.message, "{Bcast: root: 7}"); - roctxFree(&info); +TEST_F(RoctxNvmlRoctxMicrotest, Stringify_SingleEntryHasNoSeparator) { + EXPECT_EQ(Stringify("Bcast", {IntEntry("root", 7)}), "{Bcast: root: 7}"); } -TEST_F(RoctxMicrotest, Stringify_ZeroEntriesIsJustTheId) { - roctxPayloadInfo info; - roctxAlloc(&info, 1); - info.id = "Empty"; - info.numEntries = 0; - stringify(&info); - EXPECT_STREQ(info.message, "{Empty: }"); - roctxFree(&info); +TEST_F(RoctxNvmlRoctxMicrotest, Stringify_ZeroEntriesIsJustTheId) { + EXPECT_EQ(Stringify("Empty", {}), "{Empty: }"); } // Out-of-range redop/datatype values print "unknown" instead of indexing past the tables; // the highest valid values sit right next to them so an off-by-one in either bound shows. -TEST_F(RoctxMicrotest, Stringify_OutOfRangeRedOpAndDataTypePrintUnknown) { - roctxPayloadInfo info; - roctxAlloc(&info, 4); - info.id = "X"; - info.numEntries = 4; - info.payloadEntries[0] = {"lastOp", ROCTX_PAYLOAD_ENTRY_TYPE_REDOP, {}}; - info.payloadEntries[0].payload.typeRedOp = static_cast(ncclNumDevRedOps - 1); - info.payloadEntries[1] = {"badOp", ROCTX_PAYLOAD_ENTRY_TYPE_REDOP, {}}; - info.payloadEntries[1].payload.typeRedOp = static_cast(ncclNumDevRedOps); - info.payloadEntries[2] = {"lastType", ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE, {}}; - info.payloadEntries[2].payload.typeDataType = static_cast(ncclNumTypes - 1); - info.payloadEntries[3] = {"badType", ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE, {}}; - info.payloadEntries[3].payload.typeDataType = static_cast(ncclNumTypes); - stringify(&info); - EXPECT_STREQ(info.message, "{X: lastOp: SumPostDiv, badOp: unknown, lastType: b8, badType: unknown}"); - roctxFree(&info); +TEST_F(RoctxNvmlRoctxMicrotest, Stringify_OutOfRangeRedOpAndDataTypePrintUnknown) { + std::vector entries(4); + entries[0] = {"lastOp", ROCTX_PAYLOAD_ENTRY_TYPE_REDOP, {}}; + entries[0].payload.typeRedOp = static_cast(ncclNumDevRedOps - 1); + entries[1] = {"badOp", ROCTX_PAYLOAD_ENTRY_TYPE_REDOP, {}}; + entries[1].payload.typeRedOp = static_cast(ncclNumDevRedOps); + entries[2] = {"lastType", ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE, {}}; + entries[2].payload.typeDataType = static_cast(ncclNumTypes - 1); + entries[3] = {"badType", ROCTX_PAYLOAD_ENTRY_TYPE_DATATYPE, {}}; + entries[3].payload.typeDataType = static_cast(ncclNumTypes); + EXPECT_EQ(Stringify("X", entries), "{X: lastOp: SumPostDiv, badOp: unknown, lastType: b8, badType: unknown}"); } -TEST_F(RoctxMicrotest, Stringify_UnmappedRoctxTypeHitsTheDefaultBranch) { - roctxPayloadInfo info; - roctxAlloc(&info, 1); - info.id = "X"; - info.numEntries = 1; - info.payloadEntries[0] = {"weird", static_cast(ROCTX_PAYLOAD_NUM_ENTRY_TYPES), {}}; - stringify(&info); - EXPECT_STREQ(info.message, "{X: weird: unknown roctx payload type}"); - roctxFree(&info); +TEST_F(RoctxNvmlRoctxMicrotest, Stringify_UnmappedRoctxTypeHitsTheDefaultBranch) { + roctxPayloadSchemaEntryInfo weird{"weird", static_cast(ROCTX_PAYLOAD_NUM_ENTRY_TYPES), {}}; + EXPECT_EQ(Stringify("X", {weird}), "{X: weird: unknown roctx payload type}"); } -TEST_F(RoctxMicrotest, Stringify_MissingBufferIsANoOp) { +TEST_F(RoctxNvmlRoctxMicrotest, Stringify_MissingBufferIsANoOp) { roctxPayloadInfo info; // payloadEntries and message both null info.id = "X"; info.numEntries = 0; @@ -359,7 +387,7 @@ TEST_F(RoctxMicrotest, Stringify_MissingBufferIsANoOp) { // Entries present but no message buffer: still a no-op, checked separately from the // both-null case above because the guard is `!entries || !message`. -TEST_F(RoctxMicrotest, Stringify_EntriesWithoutMessageIsANoOp) { +TEST_F(RoctxNvmlRoctxMicrotest, Stringify_EntriesWithoutMessageIsANoOp) { roctxPayloadSchemaEntryInfo entry{"n", ROCTX_PAYLOAD_ENTRY_TYPE_INT, {}}; roctxPayloadInfo info; info.payloadEntries = &entry; @@ -371,7 +399,7 @@ TEST_F(RoctxMicrotest, Stringify_EntriesWithoutMessageIsANoOp) { // A mapping to a roctx type outside the enum reaches extractPayloadInfo's default branch, // which stores the type but no payload, and stringify then prints the unknown-type text. -TEST_F(RoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefaultBranch) { +TEST_F(RoctxNvmlRoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefaultBranch) { const uint64_t nvtxType = NVTX_PAYLOAD_ENTRY_TYPE_INT + 2; ASSERT_EQ(nvtxToRoctx.count(nvtxType), 0u); const auto bogus = static_cast(ROCTX_PAYLOAD_NUM_ENTRY_TYPES); @@ -382,7 +410,6 @@ TEST_F(RoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefaultBranch) { roctxPayloadInfo info; roctxAlloc(&info, 1); extractPayloadInfo(schema, &data, 1, "D", &info); - nvtxToRoctx.erase(nvtxType); EXPECT_EQ(info.payloadEntries[0].type, bogus); EXPECT_STREQ(info.message, "{D: w: unknown roctx payload type}"); roctxFree(&info); @@ -390,7 +417,7 @@ TEST_F(RoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefaultBranch) { // Pinned as observed: an nvtx type with no mapping is looked up with std::map::operator[], // which inserts it as ROCTX_PAYLOAD_ENTRY_TYPE_INT (value 0) and decodes the field as an int. -TEST_F(RoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInsertedIntoTheMap) { +TEST_F(RoctxNvmlRoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInsertedIntoTheMap) { const uint64_t unmapped = NVTX_PAYLOAD_ENTRY_TYPE_INT + 1; ASSERT_EQ(nvtxToRoctx.count(unmapped), 0u); int value = 42; @@ -401,8 +428,7 @@ TEST_F(RoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInsertedIntoTheM extractPayloadInfo(schema, &data, 1, "U", &info); EXPECT_EQ(info.payloadEntries[0].type, ROCTX_PAYLOAD_ENTRY_TYPE_INT); EXPECT_STREQ(info.message, "{U: v: 42}"); - EXPECT_EQ(nvtxToRoctx.count(unmapped), 1u); - nvtxToRoctx.erase(unmapped); // keep the global table as other tests expect it + EXPECT_EQ(nvtxToRoctx.count(unmapped), 1u); // TearDown restores the table roctxFree(&info); } @@ -410,7 +436,7 @@ TEST_F(RoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInsertedIntoTheM // roctx_scoped_range_in // --------------------------------------------------------------------------- -TEST_F(RoctxMicrotest, LogRoctxParam_IsReadFromRcclLogRoctxWithDefaultZero) { +TEST_F(RoctxNvmlRoctxMicrotest, LogRoctxParam_IsReadFromRcclLogRoctxWithDefaultZero) { std::string seenEnv; int64_t seenDefault = -1; g_loadParam = [&](const char* env, int64_t deftVal) { @@ -423,7 +449,7 @@ TEST_F(RoctxMicrotest, LogRoctxParam_IsReadFromRcclLogRoctxWithDefaultZero) { EXPECT_EQ(seenDefault, 0); } -TEST_F(RoctxMicrotest, ScopedRange_LoggingOff_DoesNothing) { +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_LoggingOff_DoesNothing) { SetLogRoctx(0); { roctx_scoped_range_in range("msg"); @@ -435,49 +461,65 @@ TEST_F(RoctxMicrotest, ScopedRange_LoggingOff_DoesNothing) { // One range is one event on enter and one on exit: a roctracer push/pop with ROCTX_ENABLE, // otherwise the WARN that stands in for each. -TEST_F(RoctxMicrotest, ScopedRange_LoggingOn_OneEventOnEnterAndOneOnExit) { +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_LoggingOn_OneEventOnEnterAndOneOnExit) { SetLogRoctx(1); { roctx_scoped_range_in range("msg"); EXPECT_EQ(RangeEvents(), 1u); ExpectPushedMessage("msg"); } - EXPECT_EQ(RangeEvents(), 2u); -#ifndef ROCTX_ENABLE - for (const auto& m : g_logged) EXPECT_NE(m.text.find("ROCTX_ENABLE is not defined"), std::string::npos); -#endif + ExpectOneRangeEnteredAndExited(); } -TEST_F(RoctxMicrotest, ScopedRange_DefaultConstructorDelegatesToTheMessageConstructor) { +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_DefaultConstructorDelegatesToTheMessageConstructor) { SetLogRoctx(1); { roctx_scoped_range_in range; ExpectPushedMessage(""); } - EXPECT_EQ(RangeEvents(), 2u); + ExpectOneRangeEnteredAndExited(); } // The payload constructor builds the message from the schema (pushed verbatim with ROCTX_ENABLE) -// and the destructor frees it in both builds; LeakSanitizer would flag a missing roctxFree. -TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOn_PushesTheStringifiedPayload) { +// and the destructor frees it in both builds; under ASan, LeakSanitizer flags a missing roctxFree. +// The two mallocs (entries, message) are checked directly too, since without ROCTX_ENABLE the +// pushed message is not observable. +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOn_PushesTheStringifiedPayload) { SetLogRoctx(1); + int* mallocs = CountMallocs(); MixedPayload payload{1, 2, ncclDevSum, ncclInt8}; nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "n", offsetof(MixedPayload, i))}; nvtxPayloadData_t data{0, sizeof(payload), &payload}; { roctx_scoped_range_in range(schema, &data, 1, "RCCL_AllReduce"); + EXPECT_EQ(*mallocs, 2); EXPECT_EQ(RangeEvents(), 1u); ExpectPushedMessage("{RCCL_AllReduce: n: 1}"); } - EXPECT_EQ(RangeEvents(), 2u); + ExpectOneRangeEnteredAndExited(); +} + +// roctxAlloc fails inside the constructor, leaving the message null: construction and +// destruction must still complete without crashing and report one enter and one exit. +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOnAndAllocFailure_StillEntersAndExits) { + SetLogRoctx(1); + int* mallocs = FailMallocAfter(0); + int value = 1; + nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "n", 0)}; + nvtxPayloadData_t data{0, sizeof(value), &value}; + { roctx_scoped_range_in range(schema, &data, 1, "RCCL_AllReduce"); } + EXPECT_EQ(*mallocs, 1); + ExpectOneRangeEnteredAndExited(); } -TEST_F(RoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOff_DoesNothing) { +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOff_DoesNothing) { SetLogRoctx(0); + int* mallocs = CountMallocs(); int value = 1; nvtxPayloadSchemaEntry_t schema[] = {Entry(NVTX_PAYLOAD_ENTRY_TYPE_INT, "n", 0)}; nvtxPayloadData_t data{0, sizeof(value), &value}; { roctx_scoped_range_in range(schema, &data, 1, "RCCL_AllReduce"); } + EXPECT_EQ(*mallocs, 0); EXPECT_EQ(RangeEvents(), 0u); EXPECT_TRUE(g_logged.empty()); } @@ -490,13 +532,13 @@ namespace { nvmlDevice_t FakeDevice() { return reinterpret_cast(0x1234); } } // namespace -TEST(NvmlStubMicrotest, LifecycleCallsSucceed) { +TEST(RoctxNvmlStubMicrotest, LifecycleCallsSucceed) { EXPECT_EQ(ncclNvmlSymbols(), ncclSuccess); EXPECT_EQ(ncclNvmlInit(), ncclSuccess); EXPECT_EQ(ncclNvmlShutdown(), ncclSuccess); } -TEST(NvmlStubMicrotest, ReportsZeroDevices) { +TEST(RoctxNvmlStubMicrotest, ReportsZeroDevices) { EXPECT_EQ(ncclNvmlDeviceCount, 0); unsigned int count = 99; EXPECT_EQ(ncclNvmlDeviceGetCount(&count), ncclSuccess); @@ -504,7 +546,7 @@ TEST(NvmlStubMicrotest, ReportsZeroDevices) { } // The successful getters write a fixed value; check it lands rather than a stale input. -TEST(NvmlStubMicrotest, SucceedingGettersWriteFixedValues) { +TEST(RoctxNvmlStubMicrotest, SucceedingGettersWriteFixedValues) { unsigned index = 99; EXPECT_EQ(ncclNvmlDeviceGetIndex(FakeDevice(), &index), ncclSuccess); EXPECT_EQ(index, 0u); @@ -520,7 +562,7 @@ TEST(NvmlStubMicrotest, SucceedingGettersWriteFixedValues) { } // The failing lookups return ncclSystemError and leave their out-params alone. -TEST(NvmlStubMicrotest, DeviceLookupsFailWithoutWritingOutputs) { +TEST(RoctxNvmlStubMicrotest, DeviceLookupsFailWithoutWritingOutputs) { nvmlDevice_t sentinel = FakeDevice(); nvmlDevice_t dev = sentinel; EXPECT_EQ(ncclNvmlDeviceGetHandleByPciBusId("0000:00:00.0", &dev), ncclSystemError); @@ -549,7 +591,7 @@ TEST(NvmlStubMicrotest, DeviceLookupsFailWithoutWritingOutputs) { // These fail too, but clear their output first so a caller that ignores the error // reads a defined value instead of stack garbage. -TEST(NvmlStubMicrotest, FailingQueriesClearTheirOutputs) { +TEST(RoctxNvmlStubMicrotest, FailingQueriesClearTheirOutputs) { char name[8] = "garbage"; EXPECT_EQ(ncclNvmlDeviceGetName(FakeDevice(), name, sizeof(name)), ncclSystemError); EXPECT_EQ(name[0], '\0'); @@ -566,7 +608,7 @@ TEST(NvmlStubMicrotest, FailingQueriesClearTheirOutputs) { } // A zero length means "no buffer": nothing may be written, not even the terminator. -TEST(NvmlStubMicrotest, ZeroLengthBuffersAreNotWritten) { +TEST(RoctxNvmlStubMicrotest, ZeroLengthBuffersAreNotWritten) { char name[2] = {'n', 'x'}; EXPECT_EQ(ncclNvmlDeviceGetName(FakeDevice(), name, 0), ncclSystemError); EXPECT_EQ(name[0], 'n'); @@ -576,7 +618,7 @@ TEST(NvmlStubMicrotest, ZeroLengthBuffersAreNotWritten) { EXPECT_EQ(version[0], 'v'); } -TEST(NvmlStubMicrotest, FieldValuesMarksEveryRequestedFieldNotSupported) { +TEST(RoctxNvmlStubMicrotest, FieldValuesMarksEveryRequestedFieldNotSupported) { nvmlFieldValue_t values[3]; std::memset(values, 0, sizeof(values)); for (auto& v : values) v.nvmlReturn = NVML_SUCCESS; diff --git a/projects/rccl/test/host/run_host_tests.sh b/projects/rccl/test/host/run_host_tests.sh index 8047a8c653d6..4591a0bf5f14 100755 --- a/projects/rccl/test/host/run_host_tests.sh +++ b/projects/rccl/test/host/run_host_tests.sh @@ -168,6 +168,9 @@ do_host_tests() { "rccl-UnitTestsMicroCe:$SCRIPT_DIR/host_tests_micro_ce.xml" "rccl-UnitTestsMicroClientSupport:$SCRIPT_DIR/host_tests_micro_client_support.xml" "rccl-UnitTestsMicroRoctxNvml:$SCRIPT_DIR/host_tests_micro_roctx_nvml.xml" + # roctx.cc gates its roctracer calls on ROCTX_ENABLE at the preprocessor; this + # variant builds that arm against a stand-in roctracer header. See test/host/CMakeLists.txt. + "rccl-UnitTestsMicroRoctxNvml-roctx:$SCRIPT_DIR/host_tests_micro_roctx_nvml_roctx.xml" "rccl-UnitTestsMicroGinHost:$SCRIPT_DIR/host_tests_micro_gin_host.xml" "rccl-UnitTestsMicroDiagnostics:$SCRIPT_DIR/host_tests_micro_diagnostics.xml" ) diff --git a/projects/rccl/test/test_categories_micro.yaml b/projects/rccl/test/test_categories_micro.yaml index 566e6d1f387a..9675aa4eb966 100644 --- a/projects/rccl/test/test_categories_micro.yaml +++ b/projects/rccl/test/test_categories_micro.yaml @@ -114,9 +114,8 @@ test_categories: # ce_coll.cc (rccl-UnitTestsMicroCe) - "CeCollMicrotest.*" - "LaunchPaths/CeAlltoAllvSyncMicrotest.*" - # misc/roctx.cc, misc/nvmlwrap_stub.cc (rccl-UnitTestsMicroRoctxNvml) - - "RoctxMicrotest.*" - - "NvmlStubMicrotest.*" + # misc/roctx.cc, misc/nvmlwrap_stub.cc (rccl-UnitTestsMicroRoctxNvml[-roctx]) + - "RoctxNvml*" # dev_runtime.cc - "Alloc*" - "Comm*" From 090feb04b69afebc677fffc4d933fb04f9922308 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Tue, 13 Oct 2026 15:21:32 -0700 Subject: [PATCH 06/16] docs(rccl): note the in-tree ROCTX_ENABLE inheritance in the RoctxNvml README entry Co-Authored-By: Claude --- projects/rccl/test/host/MICROTEST_README.md | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index 92709661e2bc..8a3573f40f06 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -257,10 +257,12 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: and `ras-diagnostics-gpu-test.cc` defines its own NVML seams; the TU supplies its own `ncclDebugLog` and `g_loadParam`, so no shared fakes are linked. `roctx.cc` gates its `roctxRangePushA`/`roctxRangePop` calls on - `ROCTX_ENABLE` at the preprocessor: the base binary covers the WARN fallback, - and the standalone-only `-roctx` variant defines `ROCTX_ENABLE` against the - stand-in `fakes/roctx_include/roctracer/roctx.h`, with the test TU recording - each push and pop. + `ROCTX_ENABLE` at the preprocessor. Standalone, the base binary covers the + WARN fallback and the standalone-only `-roctx` variant defines `ROCTX_ENABLE` + against the stand-in `fakes/roctx_include/roctracer/roctx.h`, with the test TU + recording each push and pop. In-tree, the base binary inherits `ROCTX_ENABLE` + from the `rccl` target whenever roctracer is found, so it compiles the + roctracer arm instead and no in-tree binary builds the WARN fallback. - **`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.*`, From 5083ca2bb6493feb565642f71901b3cb82f4abe4 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 11:53:48 -0700 Subject: [PATCH 07/16] test(rccl): pin cuMem off in the DMABUF_ENABLE=0 test, check the export gate stays unarmed when ROCr lacks DMA-BUF Co-Authored-By: Claude --- projects/rccl/test/host/rocmwrap-test.cc | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/rccl/test/host/rocmwrap-test.cc b/projects/rccl/test/host/rocmwrap-test.cc index 4ad99b6c8ac8..aaf066d69b78 100644 --- a/projects/rccl/test/host/rocmwrap-test.cc +++ b/projects/rccl/test/host/rocmwrap-test.cc @@ -657,6 +657,9 @@ TEST_F(RocmwrapMicrotest, Init_RecordsTheCuMemCapability) { TEST_F(RocmwrapMicrotest, Init_DmaBufDisabledByParamStopsBeforeHsaInit) { g_params["DMABUF_ENABLE"] = 0; + // Pin cuMem off: in-window builds would otherwise auto-enable it on this gfx1250 fixture, and + // ncclCuMemEnable() force-enables DMA-BUF ahead of the NCCL_DMABUF_ENABLE=0 arm under test. + g_params["CUMEM_ENABLE"] = 0; initOnceFunc(); EXPECT_EQ(initResult, ncclSystemError); EXPECT_TRUE(Logged(NCCL_LOG_INFO, "Dmabuf feature disabled")); @@ -687,6 +690,8 @@ TEST_F(RocmwrapMicrotest, Init_RocrWithoutDmaBufSupportStopsBeforeHsaInit) { EXPECT_EQ(initResult, ncclSystemError); EXPECT_TRUE(Logged(NCCL_LOG_INFO, "does not support dmabuf")); EXPECT_EQ(g_hsaInitCalls, 0); + // net.cc and coll_net.cc gate on the pointer itself, not on initResult. + EXPECT_EQ(pfn_hsa_amd_portable_export_dmabuf, nullptr); } } From 0033707060bb0fec8f66eafe4587b128cb20fd74 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 12:51:06 -0700 Subject: [PATCH 08/16] test(rccl): add host microtests for misc/amdsmi_wrap.cc Covers init/shutdown caching, GPU count, PCI bus-id lookups both ways, link info, firmware version, fabric discovery (including the amd_smi fabric ABI probe and its sysfs fallbacks) and telemetry, across the WSL2, amd_smi_lib and ARSMI backends, in a new rccl-UnitTestsMicroAmdsmiWrap binary. Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 21 + projects/rccl/test/host/MICROTEST_README.md | 15 + projects/rccl/test/host/amdsmi-wrap-test.cc | 1409 +++++++++++++++++ projects/rccl/test/host/run_host_tests.sh | 1 + .../test_categories_micro_amdsmi_wrap.yaml | 47 + .../test_runner/configs/ci-precheckin.json | 17 + .../scripts/test_runner/lib/test_executor.py | 3 +- 7 files changed, 1512 insertions(+), 1 deletion(-) create mode 100644 projects/rccl/test/host/amdsmi-wrap-test.cc create mode 100644 projects/rccl/test/test_categories_micro_amdsmi_wrap.yaml diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 1344ae8b3bf2..6ea508320e4b 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -373,6 +373,17 @@ function(rccl_define_micro_source_lists) ../common/main_altrsmi.cpp PARENT_SCOPE) + # rccl-UnitTestsMicroAmdsmiWrap -- src/misc/amdsmi_wrap.cc. Its own binary: fakes/amdsmi_fakes.cc, + # linked by most other targets, defines the amd_smi_* wrappers this target compiles for real. The + # test TU supplies ncclDebugLog, g_loadParam and the ARSMI_* (alt_rsmi.cc) stand-ins itself; + # utils.cc is the real thing for busIdToInt64. + set(TEST_MICRO_AMDSMI_WRAP_SOURCE_FILES + amdsmi-wrap-test.cc + fakes/hip_fakes.cc + ${RCCL_HOST_HIPIFY_DIR}/src/misc/utils.cc # real busIdToInt64 + ../common/main_altrsmi.cpp + PARENT_SCOPE) + # src/diagnostics/p2p.cc; fakes transport/p2p.cc, which rccl-UnitTestsMicro compiles. set(TEST_MICRO_DIAGNOSTICS_SOURCE_FILES diag-p2p-test.cc @@ -1055,6 +1066,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-UnitTestsMicroAmdsmiWrap +# --------------------------------------------------------------------------- +rccl_add_micro_binary(rccl-UnitTestsMicroAmdsmiWrap + SOURCES ${TEST_MICRO_AMDSMI_WRAP_SOURCE_FILES} + EXTRA_DEFS + ${_MICRO_MODE_DEFS} + AMDSMI_WRAP_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/amdsmi_wrap.cc" + CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_amdsmi_wrap.yaml") + # --------------------------------------------------------------------------- # rccl-UnitTestsMicroSymKernels # --------------------------------------------------------------------------- diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index 1ff07e6588f5..d5ec35230aa8 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -310,6 +310,19 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: `NCCL_GIN_PROXY_NTHREADS` progress-thread assignment. Its own binary, not sharing `rccl-UnitTestsMicro`: `gin-plugin-init-test.cc` already defines `ncclParamGinEnable` there. See `test_categories_micro_gin_host.yaml`. +- **`rccl-UnitTestsMicroAmdsmiWrap`** — `src/misc/amdsmi_wrap.cc` (via + `AMDSMI_WRAP_CC_PATH`, from `amdsmi-wrap-test.cc`); suite + `AmdsmiWrapMicrotest.*`. Covers init/shutdown and their caching, GPU count, + PCI bus-id lookups both ways, link info, firmware version, fabric discovery + (including the amd_smi fabric ABI probe and its sysfs fallbacks) and + telemetry, each across the WSL2, amd_smi_lib and ARSMI backends. Its own + binary because `fakes/amdsmi_fakes.cc`, linked by most other targets, + defines the `amd_smi_*` wrappers this file compiles for real. The amd_smi + library is a scripted `FakeSmi` installed through the file's own `pfn_*` + table; `ARSMI_*` (`alt_rsmi.cc`) is defined in the test TU, and `access`, + `dlopen`/`dlsym` and `hipDeviceGetByPCIBusId` are seamed inside + `amdsmi_wrap.cc` only; `utils.cc` is real for `busIdToInt64`. See + `test_categories_micro_amdsmi_wrap.yaml`. - **`rccl-UnitTestsMicroDiagnostics`**: `src/diagnostics/p2p.cc` (via `DIAG_P2P_CC_PATH`, suite `DiagP2pMicrotest.*`) and `src/diagnostics/ib_write_bw.cc` (via `DIAG_IB_WRITE_BW_CC_PATH`, suite `DiagIbWriteBwMicrotest.*`). Its own binary: it fakes the @@ -829,6 +842,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-UnitTestsMicroAmdsmiWrap`, `rccl-UnitTestsMicroSymKernels` and `rccl-UnitTestsMicroTaskPrep` — **without configuring/building all of librccl**. It compiles just the tests + fakes + the hipified unit-under-test sources. @@ -869,6 +883,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-UnitTestsMicroAmdsmiWrap # src/misc/amdsmi_wrap.cc ./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/p2p.cc tests ./build/rccl-UnitTestsMicroCe # src/ce_coll.cc tests ./build/rccl-HostUnitTests diff --git a/projects/rccl/test/host/amdsmi-wrap-test.cc b/projects/rccl/test/host/amdsmi-wrap-test.cc new file mode 100644 index 000000000000..c176b8aba7b9 --- /dev/null +++ b/projects/rccl/test/host/amdsmi-wrap-test.cc @@ -0,0 +1,1409 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/misc/amdsmi_wrap.cc: RCCL's wrapper over GPU +// enumeration, PCI bus IDs, link topology, firmware version and UALoE/UALink +// fabric discovery. Every query has three backends: WSL2 (HIP calls only), +// amd_smi_lib (dlopen'd, opted into with RCCL_USE_AMD_SMI_LIB=1) and the +// internal ARSMI sysfs reader (the default). +// +// The amd_smi_lib backend is reached through the file's own pfn_* table, which +// this TU can assign because it #includes the .cc; FakeSmi below stands in for +// the library as a list of GPUs. ARSMI_* (src/misc/alt_rsmi.cc) is defined +// here as scriptable stand-ins. access (the WSL2 probe), dlopen/dlsym and +// hipDeviceGetByPCIBusId are seamed inside amdsmi_wrap.cc's own bodies only; +// the other HIP calls go through fakes/hip_fakes.cc, and utils.cc's +// busIdToInt64 is the real code. +// +// The file has process-wide state the tests reset by hand (see Reset()): the +// init/fabric caches, the WSL2 flag, and the dlopen handle. The handle is a +// function-local static, so the two dlopen tests run in forked children. +// +// The fabric ABI probe compares what the loaded library writes against the +// layout this build was compiled for. FakeSmi writes the compiled layout's +// extent by default, so the typed path is taken whatever amd_smi header the +// build found; the mismatch tests write the other extents. + +#include + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "fakes/hip_fakes.h" + +#include "nccl.h" +#include "debug.h" +#include "amdsmi_wrap.h" +#include "alt_rsmi.h" +#include "core.h" +#include "utils.h" + +// --- Logging and params: this binary links no shared nccl fakes. ------------- + +uint32_t ncclDebugLevelMask = NCCL_DEBUG_LEVEL_MASK_UNINITIALIZED; // log everything +uint64_t ncclDebugMask = ~uint64_t(0); +thread_local int ncclDebugNoWarn = 0; // utils.cc reads it + +namespace { +struct LoggedMessage { + ncclDebugLogLevel level; + std::string text; +}; +std::vector g_logged; +std::map g_params; +} // namespace + +void ncclDebugLog(ncclDebugLogLevel level, unsigned long /*flags*/, const char* /*filefunc*/, int /*line*/, + const char* fmt, ...) { + char buf[512]; + va_list ap; + va_start(ap, fmt); + std::vsnprintf(buf, sizeof(buf), fmt, ap); + va_end(ap); + g_logged.push_back({level, buf}); +} + +std::function g_loadParam = [](const char* env, int64_t deftVal) { + auto it = g_params.find(env); + return it == g_params.end() ? deftVal : it->second; +}; + +#include "fakes/param_redirect.h" + +// --- ARSMI (src/misc/alt_rsmi.cc): scriptable stand-ins. ---------------------- + +namespace { +struct FakeArsmi { + int initResult = 0; + int numDevicesResult = 0; + uint32_t numDevices = 0; + std::vector bdfIds; // ARSMI_dev_pci_id_get, by index + int pciIdResult = 0; + ARSMI_linkInfo link{}; + int linkResult = 0; + uint64_t fwVersion = 0; + int fwResult = 0; + std::map fabric; // devices missing here report ENODEV + int initCalls = 0; +}; +FakeArsmi g_arsmi; +} // namespace + +int ARSMI_init(void) { + ++g_arsmi.initCalls; + return g_arsmi.initResult; +} +int ARSMI_get_num_devices(uint32_t* n) { + if (g_arsmi.numDevicesResult == 0) *n = g_arsmi.numDevices; + return g_arsmi.numDevicesResult; +} +int ARSMI_dev_pci_id_get(uint32_t i, uint64_t* id) { + if (g_arsmi.pciIdResult != 0) return g_arsmi.pciIdResult; + *id = i < g_arsmi.bdfIds.size() ? g_arsmi.bdfIds[i] : 0; + return 0; +} +int ARSMI_topo_get_link_info(uint32_t, uint32_t, ARSMI_linkInfo* info) { + *info = g_arsmi.link; + return g_arsmi.linkResult; +} +int ARSMI_get_fw_version(uint32_t, uint64_t* v) { + *v = g_arsmi.fwVersion; + return g_arsmi.fwResult; +} +int ARSMI_get_fabric_info(uint32_t i, struct ARSMI_fabricInfo* info) { + auto it = g_arsmi.fabric.find(i); + if (it == g_arsmi.fabric.end()) return 19; // ENODEV + *info = it->second; + return 0; +} +const char* ARSMI_fabric_telem_id_to_string(uint64_t) { return "ARSMI_NAME"; } + +// --- Seams scoped to amdsmi_wrap.cc's own function bodies. -------------------- + +namespace { +std::function g_access; +std::function g_dlopen; +std::map g_dlsymTable; +std::function g_hipDeviceGetByPCIBusId; +} // namespace + +static int AmdsmiTestAccess(const char* path, int mode) { return g_access(path, mode); } +static void* AmdsmiTestDlopen(const char* name, int flags) { return g_dlopen(name, flags); } +static void* AmdsmiTestDlsym(void*, const char* name) { + auto it = g_dlsymTable.find(name); + return it == g_dlsymTable.end() ? nullptr : it->second; +} +static hipError_t AmdsmiTestDeviceGetByPCIBusId(int* dev, const char* busId) { + return g_hipDeviceGetByPCIBusId(dev, busId); +} + +#define access AmdsmiTestAccess +#define dlopen AmdsmiTestDlopen +#define dlsym AmdsmiTestDlsym +#define hipDeviceGetByPCIBusId AmdsmiTestDeviceGetByPCIBusId +#include AMDSMI_WRAP_CC_PATH +#undef hipDeviceGetByPCIBusId +#undef dlsym +#undef dlopen +#undef access + +namespace { + +// Write extent of each runtime layout amdSmiDetectFabricRuntimeLayout recognizes. +constexpr size_t kExtent8Gpu = kAmdSmiFabricInfo8GpuSize; +constexpr size_t kExtent16Gpu = kAmdSmiFabricInfo16GpuSize; +constexpr size_t kExtentExtendedUnion = kAmdSmiFabricV1PayloadEnd; +// The extent a library matching this build's header writes, so the typed path is taken. +constexpr size_t kCompiledExtent = + amdSmiFabricLayoutIs8Gpu ? kExtent8Gpu : (amdSmiFabricLayoutIs16Gpu ? kExtent16Gpu : kExtentExtendedUnion); + +template +void SetFabricVersion(FabricInfoT& info, uint32_t version) { + if constexpr (amdSmiFabricInfoIsFlat::value) { + info.fabric_version = version; + } else { + info.fabric_info.version = version; + } +} + +// One amd_smi_lib processor. Handles are this struct's address, so a test can +// tell which processor a call was made with. +struct FakeProc { + bool gpu = true; + uint32_t hipId = 0; + uint64_t bdfId = 0; + uint64_t fwVersion = 0; + uint8_t numFwInfo = 1; + bool fabricOk = true; + uint32_t fabricVersion = AMDSMI_FABRIC_INFO_CURRENT_VERSION; + amdsmi_fabric_type_t fabricType = AMDSMI_FABRIC_TYPE_UALOE; + amdsmi_fabric_accelerator_vpod_state_t accelState = AMDSMI_FABRIC_ACCELERATOR_VPOD_STATE_ACTIVE; + uint32_t accelId = 0; + uint32_t bandwidth = 0; + uint32_t latency = 0; + uint32_t vpodId = 0; + uint32_t vpodSize = 0; + uint32_t ppodSize = 0; + uint8_t ppodId[16] = {}; +}; + +// The amd_smi library as a set of sockets, each a list of processors. +struct FakeSmi { + std::vector> sockets; + amdsmi_status_t initStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t versionStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t shutdownStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t socketsStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t procsStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t enumStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t bdfStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t fromBdfStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t fwStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t telemStatus = AMDSMI_STATUS_SUCCESS; + amdsmi_status_t fabricProbeStatus = AMDSMI_STATUS_SUCCESS; + size_t fabricExtent = kCompiledExtent; // bytes the library writes, from offset 0 + uint32_t libMajor = 27; + amdsmi_link_type_t linkType = AMDSMI_LINK_TYPE_PCIE; + uint64_t linkHops = 3; + uint64_t minBw = 0; + uint64_t maxBw = 0; + int initCalls = 0; + int shutdownCalls = 0; + int fabricInfoCalls = 0; + int telemAllocCalls = 0; + int telemGetCalls = 0; + int telemFreeCalls = 0; + amdsmi_bdf_t lastBdf{}; + amdsmi_processor_handle lastLinkSrc = nullptr; + amdsmi_processor_handle lastLinkDst = nullptr; + amdsmi_processor_handle lastTelemHandle = nullptr; + uint32_t lastTelemMask = 0; + const char* telemName = "OUTPARAM_NAME"; + amdsmi_status_t telemNameStatus = AMDSMI_STATUS_SUCCESS; +}; +FakeSmi g_smi; + +FakeProc* Proc(amdsmi_processor_handle h) { return static_cast(h); } + +amdsmi_status_t FakeInit(uint64_t) { + ++g_smi.initCalls; + return g_smi.initStatus; +} +amdsmi_status_t FakeShutDown() { + ++g_smi.shutdownCalls; + return g_smi.shutdownStatus; +} +amdsmi_status_t FakeStatusString(amdsmi_status_t, const char** s) { + *s = "FAKE_STATUS"; + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeLibVersion(amdsmi_version_t* v) { + if (g_smi.versionStatus != AMDSMI_STATUS_SUCCESS) return g_smi.versionStatus; + std::memset(v, 0, sizeof(*v)); + v->major = g_smi.libMajor; + v->minor = 1; + v->release = 2; + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeSocketHandles(uint32_t* count, amdsmi_socket_handle* handles) { + if (g_smi.socketsStatus != AMDSMI_STATUS_SUCCESS) return g_smi.socketsStatus; + if (handles != nullptr) { + for (uint32_t i = 0; i < *count && i < g_smi.sockets.size(); ++i) handles[i] = &g_smi.sockets[i]; + } + *count = static_cast(g_smi.sockets.size()); + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeProcessorHandles(amdsmi_socket_handle socket, uint32_t* count, amdsmi_processor_handle* handles) { + if (g_smi.procsStatus != AMDSMI_STATUS_SUCCESS) return g_smi.procsStatus; + auto& procs = *static_cast*>(socket); + if (handles != nullptr) { + for (uint32_t i = 0; i < *count && i < procs.size(); ++i) handles[i] = &procs[i]; + } + *count = static_cast(procs.size()); + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeProcessorType(amdsmi_processor_handle h, processor_type_t* type) { + *type = Proc(h)->gpu ? AMDSMI_PROCESSOR_TYPE_AMD_GPU : AMDSMI_PROCESSOR_TYPE_AMD_CPU; + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeEnumerationInfo(amdsmi_processor_handle h, amdsmi_enumeration_info_t* info) { + if (g_smi.enumStatus != AMDSMI_STATUS_SUCCESS) return g_smi.enumStatus; + std::memset(info, 0, sizeof(*info)); + info->hip_id = Proc(h)->hipId; + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeBdfId(amdsmi_processor_handle h, uint64_t* id) { + if (g_smi.bdfStatus != AMDSMI_STATUS_SUCCESS) return g_smi.bdfStatus; + *id = Proc(h)->bdfId; + return AMDSMI_STATUS_SUCCESS; +} +// Looks the BDF up the way the library does: the first processor whose bdfId matches. +amdsmi_status_t FakeHandleFromBdf(amdsmi_bdf_t bdf, amdsmi_processor_handle* h) { + g_smi.lastBdf = bdf; + if (g_smi.fromBdfStatus != AMDSMI_STATUS_SUCCESS) return g_smi.fromBdfStatus; + for (auto& s : g_smi.sockets) { + for (auto& p : s) { + if (p.bdfId == bdf.as_uint) { + *h = &p; + return AMDSMI_STATUS_SUCCESS; + } + } + } + return AMDSMI_STATUS_NOT_FOUND; +} +amdsmi_status_t FakeLinkType(amdsmi_processor_handle src, amdsmi_processor_handle dst, uint64_t* hops, + amdsmi_link_type_t* type) { + g_smi.lastLinkSrc = src; + g_smi.lastLinkDst = dst; + *hops = g_smi.linkHops; + *type = g_smi.linkType; + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeLinkWeight(amdsmi_processor_handle, amdsmi_processor_handle, uint64_t* weight) { + *weight = 0; + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeMinMaxBw(amdsmi_processor_handle, amdsmi_processor_handle, uint64_t* minBw, uint64_t* maxBw) { + *minBw = g_smi.minBw; + *maxBw = g_smi.maxBw; + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeFwInfo(amdsmi_processor_handle h, amdsmi_fw_info_t* info) { + if (g_smi.fwStatus != AMDSMI_STATUS_SUCCESS) return g_smi.fwStatus; + info->num_fw_info = Proc(h)->numFwInfo; + info->fw_info_list[0].fw_version = Proc(h)->fwVersion; + return AMDSMI_STATUS_SUCCESS; +} +// Writes like a whole-object library assignment: the first fabricExtent bytes, zeroed, then the v1 +// payload. Reserved bytes stay zero, which is what the layout probe keys on. +amdsmi_status_t FakeFabricInfo(amdsmi_processor_handle h, amdsmi_fabric_info_t* info) { + ++g_smi.fabricInfoCalls; + const FakeProc& p = *Proc(h); + if (!p.fabricOk) return AMDSMI_STATUS_NOT_SUPPORTED; + if (g_smi.fabricInfoCalls == 1 && g_smi.fabricProbeStatus != AMDSMI_STATUS_SUCCESS) { + return g_smi.fabricProbeStatus; + } + std::memset(info, 0, std::min(g_smi.fabricExtent, sizeof(amdSmiFabricInfoBuffer))); + SetFabricVersion(*info, p.fabricVersion); + auto* v1 = const_cast>*>( + amdSmiFabricInfoV1(*info)); + v1->fabric_type = p.fabricType; + v1->accel_state = p.accelState; + v1->accelerator_id = p.accelId; + v1->bandwidth = p.bandwidth; + v1->latency = p.latency; + v1->vpod_id = p.vpodId; + v1->vpod_size = p.vpodSize; + v1->ppod_size = p.ppodSize; + std::memcpy(v1->ppod_id, p.ppodId, sizeof(v1->ppod_id)); + return AMDSMI_STATUS_SUCCESS; +} +amdsmi_status_t FakeAllocTelem(amdsmi_processor_handle h, uint32_t mask, amdsmi_fabric_telemetry_t** t) { + ++g_smi.telemAllocCalls; + g_smi.lastTelemHandle = h; + g_smi.lastTelemMask = mask; + *t = reinterpret_cast(0x1000); + return g_smi.telemStatus; +} +amdsmi_status_t FakeGetTelem(amdsmi_processor_handle h, amdsmi_fabric_telemetry_t*) { + ++g_smi.telemGetCalls; + g_smi.lastTelemHandle = h; + return g_smi.telemStatus; +} +amdsmi_status_t FakeFreeTelem(amdsmi_processor_handle h, amdsmi_fabric_telemetry_t*) { + ++g_smi.telemFreeCalls; + g_smi.lastTelemHandle = h; + return g_smi.telemStatus; +} +amdsmi_status_t FakeTelemIdToString(uint64_t, const char** name) { + if (g_smi.telemNameStatus != AMDSMI_STATUS_SUCCESS) return g_smi.telemNameStatus; + *name = g_smi.telemName; + return AMDSMI_STATUS_SUCCESS; +} +const char* g_legacyTelemName = "LEGACY_NAME"; +const char* FakeLegacyTelemIdToString(uint64_t) { return g_legacyTelemName; } + +void InstallFakeSmiTable() { + pfn_amdsmi_init = FakeInit; + pfn_amdsmi_shut_down = FakeShutDown; + pfn_amdsmi_status_code_to_string = FakeStatusString; + pfn_amdsmi_get_lib_version = FakeLibVersion; + pfn_amdsmi_get_socket_handles = FakeSocketHandles; + pfn_amdsmi_get_processor_handles = FakeProcessorHandles; + pfn_amdsmi_get_processor_type = FakeProcessorType; + pfn_amdsmi_get_processor_handle_from_bdf = FakeHandleFromBdf; + pfn_amdsmi_get_gpu_enumeration_info = FakeEnumerationInfo; + pfn_amdsmi_get_gpu_bdf_id = FakeBdfId; + pfn_amdsmi_topo_get_link_type = FakeLinkType; + pfn_amdsmi_topo_get_link_weight = FakeLinkWeight; + pfn_amdsmi_get_minmax_bandwidth_between_processors = FakeMinMaxBw; + pfn_amdsmi_get_gpu_fabric_info = FakeFabricInfo; + pfn_amdsmi_alloc_fabric_telemetry = FakeAllocTelem; + pfn_amdsmi_get_fabric_telemetry_data = FakeGetTelem; + pfn_amdsmi_free_fabric_telemetry = FakeFreeTelem; + pfn_amdsmi_fabric_telem_id_to_string = FakeTelemIdToString; + pfn_amdsmi_get_fw_info = FakeFwInfo; +} + +void ClearSmiTable() { + pfn_amdsmi_init = nullptr; + pfn_amdsmi_shut_down = nullptr; + pfn_amdsmi_status_code_to_string = nullptr; + pfn_amdsmi_get_lib_version = nullptr; + pfn_amdsmi_get_socket_handles = nullptr; + pfn_amdsmi_get_processor_handles = nullptr; + pfn_amdsmi_get_processor_type = nullptr; + pfn_amdsmi_get_processor_handle_from_bdf = nullptr; + pfn_amdsmi_get_gpu_enumeration_info = nullptr; + pfn_amdsmi_get_gpu_bdf_id = nullptr; + pfn_amdsmi_topo_get_link_type = nullptr; + pfn_amdsmi_topo_get_link_weight = nullptr; + pfn_amdsmi_get_minmax_bandwidth_between_processors = nullptr; + pfn_amdsmi_get_gpu_fabric_info = nullptr; + pfn_amdsmi_alloc_fabric_telemetry = nullptr; + pfn_amdsmi_get_fabric_telemetry_data = nullptr; + pfn_amdsmi_free_fabric_telemetry = nullptr; + pfn_amdsmi_fabric_telem_id_to_string = nullptr; + pfn_amdsmi_get_fw_info = nullptr; +} + +class AmdsmiWrapMicrotest : public ::testing::Test { +protected: + void SetUp() override { Reset(); } + void TearDown() override { + ResetHipFakes(); + ClearSmiTable(); + } + + // Every seam back to a bare-metal host on the default ARSMI backend, and amdsmi_wrap.cc's + // process-wide state back to "never initialized". + void Reset() { + ResetHipFakes(); + g_logged.clear(); + g_params.clear(); + g_arsmi = FakeArsmi{}; + g_smi = FakeSmi{}; + g_access = [](const char*, int) { return -1; }; // no /dev/dxg: not WSL2 + g_dlopen = [](const char*, int) -> void* { return nullptr; }; + g_dlsymTable.clear(); + g_hipDeviceGetByPCIBusId = [](int*, const char*) { return hipErrorInvalidValue; }; + InstallFakeSmiTable(); + + is_wsl2 = -1; + amdSmiInitCalled = false; + amdSmiInitResult = ncclSuccess; + amdSmiLibInitialized = false; + amdSmiLibMajor.store(kAmdSmiLibVersionUnknown); + fabricInitialized = false; + threadFabricInitialized = false; + fabricInitResult = ncclSuccess; + amdsmiFabricDeviceCount = 0; + std::memset(amdsmiFabricDevices, 0, sizeof(amdsmiFabricDevices)); + } + + static void UseLib() { g_params["RCCL_USE_AMD_SMI_LIB"] = 1; } + static void UseWsl2() { + g_access = [](const char* path, int) { return std::strcmp(path, "/dev/dxg") == 0 ? 0 : -1; }; + } + + // Runs the one-time init so is_wsl2 and the backend are settled, then clears the log. + static void InitFor(bool lib) { + if (lib) UseLib(); + ASSERT_EQ(amd_smi_init(), ncclSuccess); + g_logged.clear(); + } + + // One socket with GPUs whose hip ids are 0..n-1, plus a CPU first so lookups must skip it. + static void TwoGpusBehindACpu() { + FakeProc cpu; + cpu.gpu = false; + cpu.hipId = 1; // would match hip id 1 if the type check were skipped + FakeProc g0, g1; + g0.hipId = 0; + g0.bdfId = 0x1234000000c1ull; + g1.hipId = 1; + g1.bdfId = 0x1234000000d2ull; + g_smi.sockets = {{cpu, g0, g1}}; + } + + static bool Logged(ncclDebugLogLevel level, const char* needle) { + for (const auto& m : g_logged) { + if (m.level == level && m.text.find(needle) != std::string::npos) return true; + } + return false; + } +}; + +} // namespace + +// --------------------------------------------------------------------------- +// amd_smi_init / amd_smi_shutdown +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, Init_Wsl2SkipsEveryBackend) { + UseWsl2(); + UseLib(); + EXPECT_EQ(amd_smi_init(), ncclSuccess); + EXPECT_EQ(is_wsl2, 1); + EXPECT_EQ(g_smi.initCalls, 0); + EXPECT_EQ(g_arsmi.initCalls, 0); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "WSL2 environment detected")); +} + +TEST_F(AmdsmiWrapMicrotest, Init_DefaultBackendIsArsmi) { + EXPECT_EQ(amd_smi_init(), ncclSuccess); + EXPECT_EQ(is_wsl2, 0); + EXPECT_EQ(g_arsmi.initCalls, 1); + EXPECT_EQ(g_smi.initCalls, 0); +} + +TEST_F(AmdsmiWrapMicrotest, Init_ArsmiFailureIsInternalError) { + g_arsmi.initResult = 5; + EXPECT_EQ(amd_smi_init(), ncclInternalError); + EXPECT_TRUE(Logged(NCCL_LOG_ERROR, "ARSMI failure: 5")); +} + +TEST_F(AmdsmiWrapMicrotest, Init_LibInitsAndRecordsTheLibraryMajor) { + UseLib(); + g_smi.libMajor = 26; + EXPECT_EQ(amd_smi_init(), ncclSuccess); + EXPECT_EQ(g_smi.initCalls, 1); + EXPECT_TRUE(amdSmiLibInitialized); + EXPECT_EQ(amdSmiLibMajor.load(), 26u); + EXPECT_EQ(g_arsmi.initCalls, 0); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "amdsmi_lib: version 26.1.2")); +} + +TEST_F(AmdsmiWrapMicrotest, Init_LibInitFailureLeavesTheLibraryUninitialized) { + UseLib(); + g_smi.initStatus = AMDSMI_STATUS_INVAL; + EXPECT_EQ(amd_smi_init(), ncclInternalError); + EXPECT_FALSE(amdSmiLibInitialized); + EXPECT_TRUE(Logged(NCCL_LOG_ERROR, "AMD SMI failure: FAKE_STATUS")); +} + +// The library is up even though the version query failed, so shutdown must still shut it down. +TEST_F(AmdsmiWrapMicrotest, Init_VersionFailureAfterInitStillCountsAsInitialized) { + UseLib(); + g_smi.versionStatus = AMDSMI_STATUS_INVAL; + EXPECT_EQ(amd_smi_init(), ncclInternalError); + EXPECT_TRUE(amdSmiLibInitialized); + EXPECT_EQ(amdSmiLibMajor.load(), kAmdSmiLibVersionUnknown); + EXPECT_EQ(amd_smi_shutdown(), ncclSuccess); + EXPECT_EQ(g_smi.shutdownCalls, 1); +} + +// The result is cached, failures included, until shutdown resets it. +TEST_F(AmdsmiWrapMicrotest, Init_CachesItsResultUntilShutdown) { + g_arsmi.initResult = 1; + EXPECT_EQ(amd_smi_init(), ncclInternalError); + g_arsmi.initResult = 0; + EXPECT_EQ(amd_smi_init(), ncclInternalError); + EXPECT_EQ(g_arsmi.initCalls, 1); + + EXPECT_EQ(amd_smi_shutdown(), ncclSuccess); + EXPECT_EQ(amd_smi_init(), ncclSuccess); + EXPECT_EQ(g_arsmi.initCalls, 2); +} + +TEST_F(AmdsmiWrapMicrotest, Shutdown_WithoutInitDoesNothing) { + EXPECT_EQ(amd_smi_shutdown(), ncclSuccess); + EXPECT_EQ(g_smi.shutdownCalls, 0); +} + +TEST_F(AmdsmiWrapMicrotest, Shutdown_ShutsTheLibraryDownOnceAndForgetsItsMajor) { + UseLib(); + ASSERT_EQ(amd_smi_init(), ncclSuccess); + EXPECT_EQ(amd_smi_shutdown(), ncclSuccess); + EXPECT_EQ(g_smi.shutdownCalls, 1); + EXPECT_FALSE(amdSmiLibInitialized); + EXPECT_EQ(amdSmiLibMajor.load(), kAmdSmiLibVersionUnknown); + EXPECT_EQ(amd_smi_shutdown(), ncclSuccess); + EXPECT_EQ(g_smi.shutdownCalls, 1); +} + +// The ARSMI backend has nothing to shut down. +TEST_F(AmdsmiWrapMicrotest, Shutdown_ArsmiBackendNeverCallsTheLibrary) { + ASSERT_EQ(amd_smi_init(), ncclSuccess); + EXPECT_EQ(amd_smi_shutdown(), ncclSuccess); + EXPECT_EQ(g_smi.shutdownCalls, 0); +} + +TEST_F(AmdsmiWrapMicrotest, Shutdown_LibraryFailureIsReturned) { + UseLib(); + ASSERT_EQ(amd_smi_init(), ncclSuccess); + g_smi.shutdownStatus = AMDSMI_STATUS_INVAL; + EXPECT_EQ(amd_smi_shutdown(), ncclInternalError); +} + +// dlopen runs once per process (function-local static handle), so these two run in forked children. +// The child bodies are plain functions: their lambdas' commas would split EXPECT_EXIT's arguments. +namespace { +int DlopenFailureChild() { + ClearSmiTable(); + g_params["RCCL_USE_AMD_SMI_LIB"] = 1; + static std::string opened; + g_dlopen = [](const char* name, int) -> void* { + opened = name; + return nullptr; + }; + ncclResult_t r = amd_smi_init(); + bool warned = false; + for (const auto& m : g_logged) { + warned |= m.level == NCCL_LOG_WARN && m.text.find("Failed to open") != std::string::npos; + } + return r == ncclInternalError && opened.rfind("libamd_smi.so", 0) == 0 && warned ? 0 : 1; +} + +int DlopenResolvesChild() { + ClearSmiTable(); + g_params["RCCL_USE_AMD_SMI_LIB"] = 1; + static int handle; + g_dlopen = [](const char*, int) -> void* { return &handle; }; + g_dlsymTable = { + {"amdsmi_init", reinterpret_cast(FakeInit)}, + {"amdsmi_shut_down", reinterpret_cast(FakeShutDown)}, + {"amdsmi_status_code_to_string", reinterpret_cast(FakeStatusString)}, + {"amdsmi_get_lib_version", reinterpret_cast(FakeLibVersion)}, + {"amdsmi_get_socket_handles", reinterpret_cast(FakeSocketHandles)}, + {"amdsmi_get_processor_handles", reinterpret_cast(FakeProcessorHandles)}, + {"amdsmi_get_processor_type", reinterpret_cast(FakeProcessorType)}, + {"amdsmi_get_processor_handle_from_bdf", reinterpret_cast(FakeHandleFromBdf)}, + {"amdsmi_get_gpu_enumeration_info", reinterpret_cast(FakeEnumerationInfo)}, + {"amdsmi_get_gpu_bdf_id", reinterpret_cast(FakeBdfId)}, + {"amdsmi_topo_get_link_type", reinterpret_cast(FakeLinkType)}, + {"amdsmi_topo_get_link_weight", reinterpret_cast(FakeLinkWeight)}, + {"amdsmi_get_minmax_bandwidth_between_processors", reinterpret_cast(FakeMinMaxBw)}, + {"amdsmi_get_gpu_fabric_info", reinterpret_cast(FakeFabricInfo)}, + {"amdsmi_fabric_telem_id_to_string", reinterpret_cast(FakeTelemIdToString)}, + {"amdsmi_alloc_fabric_telemetry", reinterpret_cast(FakeAllocTelem)}, + {"amdsmi_get_fabric_telemetry_data", reinterpret_cast(FakeGetTelem)}, + {"amdsmi_free_fabric_telemetry", reinterpret_cast(FakeFreeTelem)}, + {"amdsmi_get_fw_info", reinterpret_cast(FakeFwInfo)}, + }; + bool ok = amd_smi_init() == ncclSuccess && g_smi.initCalls == 1; + ok = ok && pfn_amdsmi_get_gpu_fabric_info == FakeFabricInfo && pfn_amdsmi_get_fw_info == FakeFwInfo && + pfn_amdsmi_get_minmax_bandwidth_between_processors == FakeMinMaxBw && + pfn_amdsmi_free_fabric_telemetry == FakeFreeTelem && pfn_amdsmi_get_lib_version == FakeLibVersion; + return ok ? 0 : 1; +} +} // namespace + +TEST_F(AmdsmiWrapMicrotest, Init_DlopenFailureIsInternalError) { + EXPECT_EXIT(std::exit(DlopenFailureChild()), ::testing::ExitedWithCode(0), ""); +} + +TEST_F(AmdsmiWrapMicrotest, Init_DlopenResolvesEveryEntryPointByName) { + EXPECT_EXIT(std::exit(DlopenResolvesChild()), ::testing::ExitedWithCode(0), ""); +} + +// --------------------------------------------------------------------------- +// getProcessorHandle (via amd_smi_getFirmwareVersion, its only plain caller) +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_IsAmdSmiLibOnly) { + amdsmi_processor_handle h = nullptr; + EXPECT_EQ(getProcessorHandle(0, &h), ncclSystemError); +} + +TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_MatchesTheHipIdAndSkipsNonGpus) { + UseLib(); + TwoGpusBehindACpu(); + amdsmi_processor_handle h = nullptr; + ASSERT_EQ(getProcessorHandle(1, &h), ncclSuccess); + EXPECT_EQ(h, &g_smi.sockets[0][2]); + EXPECT_EQ(getProcessorHandle(7, &h), ncclInternalError); +} + +TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_SearchesEverySocket) { + UseLib(); + FakeProc a, b; + a.hipId = 0; + b.hipId = 1; + g_smi.sockets = {{a}, {b}}; + amdsmi_processor_handle h = nullptr; + ASSERT_EQ(getProcessorHandle(1, &h), ncclSuccess); + EXPECT_EQ(h, &g_smi.sockets[1][0]); +} + +TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_LibraryFailuresAreInternalError) { + UseLib(); + TwoGpusBehindACpu(); + amdsmi_processor_handle h = nullptr; + g_smi.socketsStatus = AMDSMI_STATUS_INVAL; + EXPECT_EQ(getProcessorHandle(0, &h), ncclInternalError); + g_smi.socketsStatus = AMDSMI_STATUS_SUCCESS; + g_smi.enumStatus = AMDSMI_STATUS_INVAL; + EXPECT_EQ(getProcessorHandle(0, &h), ncclInternalError); +} + +#if !AMDSMI_DIRECT +// Built without the amd_smi header, an unresolved entry point is a quiet failure, not a crash. +TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_MissingSymbolFailsQuietly) { + UseLib(); + TwoGpusBehindACpu(); + pfn_amdsmi_get_socket_handles = nullptr; + amdsmi_processor_handle h = nullptr; + EXPECT_EQ(getProcessorHandle(0, &h), ncclInternalError); + EXPECT_FALSE(Logged(NCCL_LOG_ERROR, "AMD SMI failure")); +} +#endif + +// --------------------------------------------------------------------------- +// amd_smi_getNumDevice +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, NumDevice_Wsl2AsksHip) { + UseWsl2(); + InitFor(/*lib=*/false); + g_hipGetDeviceCount = [](int* n) { + *n = 3; + return hipSuccess; + }; + uint32_t n = 0; + EXPECT_EQ(amd_smi_getNumDevice(&n), ncclSuccess); + EXPECT_EQ(n, 3u); +} + +// Counts every processor handle in every socket, CPUs included (pinned as observed). +TEST_F(AmdsmiWrapMicrotest, NumDevice_LibCountsProcessorsAcrossSockets) { + InitFor(/*lib=*/true); + g_smi.sockets = {{FakeProc{}, FakeProc{}}, {FakeProc{}}}; + g_smi.sockets[1][0].gpu = false; + uint32_t n = 0; + EXPECT_EQ(amd_smi_getNumDevice(&n), ncclSuccess); + EXPECT_EQ(n, 3u); +} + +TEST_F(AmdsmiWrapMicrotest, NumDevice_LibFailureIsInternalError) { + InitFor(/*lib=*/true); + g_smi.sockets = {{FakeProc{}}}; + g_smi.procsStatus = AMDSMI_STATUS_INVAL; + uint32_t n = 0; + EXPECT_EQ(amd_smi_getNumDevice(&n), ncclInternalError); +} + +TEST_F(AmdsmiWrapMicrotest, NumDevice_ArsmiReportsItsCount) { + InitFor(/*lib=*/false); + g_arsmi.numDevices = 8; + uint32_t n = 0; + EXPECT_EQ(amd_smi_getNumDevice(&n), ncclSuccess); + EXPECT_EQ(n, 8u); + g_arsmi.numDevicesResult = 2; + EXPECT_EQ(amd_smi_getNumDevice(&n), ncclInternalError); +} + +// --------------------------------------------------------------------------- +// amd_smi_getDevicePciBusIdString +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, PciBusIdString_Wsl2AsksHip) { + UseWsl2(); + InitFor(/*lib=*/false); + g_hipDeviceGetPCIBusId = [](char* busId, int len, int dev) { + std::snprintf(busId, len, "wsl-%d", dev); + return hipSuccess; + }; + char busId[32] = {}; + EXPECT_EQ(amd_smi_getDevicePciBusIdString(2, busId, sizeof(busId)), ncclSuccess); + EXPECT_STREQ(busId, "wsl-2"); +} + +// Domain from bits [63:32], bus [15:8], device [7:3], function [2:0] (pinned as observed: the +// comment in the source describes the domain as [63:16]). +TEST_F(AmdsmiWrapMicrotest, PciBusIdString_ArsmiFormatsTheBdf) { + InitFor(/*lib=*/false); + g_arsmi.bdfIds = {0, 0x0000002A0000C3DBull}; + char busId[32] = {}; + EXPECT_EQ(amd_smi_getDevicePciBusIdString(1, busId, sizeof(busId)), ncclSuccess); + EXPECT_STREQ(busId, "002a:c3:1b.3"); + g_arsmi.pciIdResult = 1; + EXPECT_EQ(amd_smi_getDevicePciBusIdString(1, busId, sizeof(busId)), ncclInternalError); +} + +TEST_F(AmdsmiWrapMicrotest, PciBusIdString_LibFindsTheGpuByHipId) { + InitFor(/*lib=*/true); + TwoGpusBehindACpu(); + char busId[32] = {}; + ASSERT_EQ(amd_smi_getDevicePciBusIdString(1, busId, sizeof(busId)), ncclSuccess); + EXPECT_STREQ(busId, "1234:00:1a.2"); // bdfId 0x1234000000d2 of the GPU, not the CPU +} + +TEST_F(AmdsmiWrapMicrotest, PciBusIdString_LibDeviceNotFoundIsInternalError) { + InitFor(/*lib=*/true); + TwoGpusBehindACpu(); + char busId[32] = {}; + EXPECT_EQ(amd_smi_getDevicePciBusIdString(5, busId, sizeof(busId)), ncclInternalError); + EXPECT_TRUE(Logged(NCCL_LOG_ERROR, "device index 5 not found")); +} + +TEST_F(AmdsmiWrapMicrotest, PciBusIdString_LibBdfQueryFailureIsInternalError) { + InitFor(/*lib=*/true); + TwoGpusBehindACpu(); + g_smi.bdfStatus = AMDSMI_STATUS_INVAL; + char busId[32] = {}; + EXPECT_EQ(amd_smi_getDevicePciBusIdString(0, busId, sizeof(busId)), ncclInternalError); +} + +// The result is truncated to the caller's buffer. +TEST_F(AmdsmiWrapMicrotest, PciBusIdString_HonorsTheBufferLength) { + InitFor(/*lib=*/false); + g_arsmi.bdfIds = {0x0000002A0000C3DBull}; + char busId[6]; + std::memset(busId, 'x', sizeof(busId)); + EXPECT_EQ(amd_smi_getDevicePciBusIdString(0, busId, sizeof(busId)), ncclSuccess); + EXPECT_STREQ(busId, "002a:"); +} + +// --------------------------------------------------------------------------- +// amd_smi_getDeviceIndexByPciBusId +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_Wsl2AsksHip) { + UseWsl2(); + InitFor(/*lib=*/false); + std::string seen; + g_hipDeviceGetByPCIBusId = [&seen](int* dev, const char* busId) { + seen = busId; + *dev = 4; + return hipSuccess; + }; + uint32_t idx = 0; + EXPECT_EQ(amd_smi_getDeviceIndexByPciBusId("0000:c3:00.0", &idx), ncclSuccess); + EXPECT_EQ(idx, 4u); + EXPECT_EQ(seen, "0000:c3:00.0"); +} + +// ARSMI: the NCCL bus id is converted to the sysfs BDF encoding and matched against every device. +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_ArsmiMatchesTheConvertedBdf) { + InitFor(/*lib=*/false); + // "0000:c3:1b.3" -> busIdToInt64 0xc31b3 -> ((0)<<12) + (0xc3000>>4) + (0x1b0>>1) + 3 = 0xc300 + 0xd8 + 3 + g_arsmi.numDevices = 3; + g_arsmi.bdfIds = {0x1111, 0xc3db, 0x2222}; + uint32_t idx = 99; + EXPECT_EQ(amd_smi_getDeviceIndexByPciBusId("0000:c3:1b.3", &idx), ncclSuccess); + EXPECT_EQ(idx, 1u); +} + +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_ArsmiNoMatchIsInternalError) { + InitFor(/*lib=*/false); + g_arsmi.numDevices = 2; + g_arsmi.bdfIds = {0x1111, 0x2222}; + uint32_t idx = 99; + EXPECT_EQ(amd_smi_getDeviceIndexByPciBusId("0000:c3:1b.3", &idx), ncclInternalError); + EXPECT_EQ(idx, 99u); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "device index not found")); +} + +// The lib backend splits NCCL's bus-id integer into amd_smi's BDF fields itself. +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_LibPassesTheBdfFieldsAndReturnsTheHipId) { + InitFor(/*lib=*/true); + FakeProc gpu; + gpu.hipId = 6; + g_smi.sockets = {{gpu}}; + // busIdToInt64("002a:c3:0b.3") = 0x2ac30b3 -> function 3, device 0x0b, bus 0xc3, domain 0x2a + amdsmi_bdf_t expected{}; + expected.function_number = 3; + expected.device_number = 0x0b; + expected.bus_number = 0xc3; + expected.domain_number = 0x2a; + g_smi.sockets[0][0].bdfId = expected.as_uint; + uint32_t idx = 0; + ASSERT_EQ(amd_smi_getDeviceIndexByPciBusId("002a:c3:0b.3", &idx), ncclSuccess); + EXPECT_EQ(idx, 6u); + EXPECT_EQ(g_smi.lastBdf.function_number, 3u); + EXPECT_EQ(g_smi.lastBdf.device_number, 0x0bu); + EXPECT_EQ(g_smi.lastBdf.bus_number, 0xc3u); + EXPECT_EQ(g_smi.lastBdf.domain_number, 0x2au); +} + +// Pinned as observed, not endorsed: the device number is taken as `(busid & 0xff) >> 4`, which keeps +// only its low hex digit, so device 0x1b reaches amd_smi as 0x0b. Devices below 0x10 are unaffected. +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_LibDropsTheDeviceNumbersHighDigit) { + InitFor(/*lib=*/true); + g_smi.fromBdfStatus = AMDSMI_STATUS_NOT_FOUND; + uint32_t idx = 0; + EXPECT_EQ(amd_smi_getDeviceIndexByPciBusId("002a:c3:1b.3", &idx), ncclInternalError); + EXPECT_EQ(g_smi.lastBdf.device_number, 0x0bu); + EXPECT_EQ(g_smi.lastBdf.bus_number, 0xc3u); +} + +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_LibNonGpuIsInternalError) { + InitFor(/*lib=*/true); + FakeProc cpu; + cpu.gpu = false; + g_smi.sockets = {{cpu}}; + amdsmi_bdf_t bdf{}; + bdf.bus_number = 0xc3; + g_smi.sockets[0][0].bdfId = bdf.as_uint; + uint32_t idx = 0; + EXPECT_EQ(amd_smi_getDeviceIndexByPciBusId("0000:c3:00.0", &idx), ncclInternalError); + EXPECT_TRUE(Logged(NCCL_LOG_ERROR, "device index not found")); +} + +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_LibLookupFailureIsInternalError) { + InitFor(/*lib=*/true); + g_smi.fromBdfStatus = AMDSMI_STATUS_NOT_FOUND; + uint32_t idx = 0; + EXPECT_EQ(amd_smi_getDeviceIndexByPciBusId("0000:c3:00.0", &idx), ncclInternalError); +} + +// --------------------------------------------------------------------------- +// amd_smi_getLinkInfo +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, LinkInfo_Wsl2IsOnePcieHop) { + UseWsl2(); + InitFor(/*lib=*/false); + amdsmi_link_type_t type = AMDSMI_LINK_TYPE_XGMI; + int hops = 0, count = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(type, AMDSMI_LINK_TYPE_PCIE); + EXPECT_EQ(hops, 1); + EXPECT_EQ(count, 1); +} + +// ARSMI's link-type enum is not value-compatible with amdsmi's, so each is mapped explicitly. +TEST_F(AmdsmiWrapMicrotest, LinkInfo_ArsmiMapsTheLinkType) { + InitFor(/*lib=*/false); + struct Case { + ARSMI_IO_LINK_TYPE in; + amdsmi_link_type_t out; + }; + for (const Case& c : {Case{ARSMI_IOLINK_TYPE_PCIEXPRESS, AMDSMI_LINK_TYPE_PCIE}, + Case{ARSMI_IOLINK_TYPE_XGMI, AMDSMI_LINK_TYPE_XGMI}, + Case{ARSMI_IOLINK_TYPE_UNDEFINED, AMDSMI_LINK_TYPE_UNKNOWN}}) { + SCOPED_TRACE(static_cast(c.in)); + g_arsmi.link = ARSMI_linkInfo{}; + g_arsmi.link.type = c.in; + amdsmi_link_type_t type = AMDSMI_LINK_TYPE_NOT_APPLICABLE; + int hops = 0, count = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(type, c.out); + EXPECT_EQ(hops, 2); + EXPECT_EQ(count, 1); + } +} + +// An XGMI link with one of the known direct-link weights is one hop, its count the bandwidth ratio. +TEST_F(AmdsmiWrapMicrotest, LinkInfo_ArsmiDirectXgmiWeightsAreOneHop) { + InitFor(/*lib=*/false); + for (uint64_t weight : {15ull, 41ull, 13ull}) { + SCOPED_TRACE(weight); + g_arsmi.link = ARSMI_linkInfo{}; + g_arsmi.link.type = ARSMI_IOLINK_TYPE_XGMI; + g_arsmi.link.weight = weight; + g_arsmi.link.min_bandwidth = 50; + g_arsmi.link.max_bandwidth = 200; + amdsmi_link_type_t type; + int hops = 0, count = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(hops, 1); + EXPECT_EQ(count, 4); + } + // Any other weight is two hops, count 1; and a zero bandwidth leaves count at 1. + g_arsmi.link.weight = 16; + amdsmi_link_type_t type; + int hops = 0, count = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(hops, 2); + EXPECT_EQ(count, 1); + g_arsmi.link.weight = 15; + g_arsmi.link.min_bandwidth = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(hops, 1); + EXPECT_EQ(count, 1); +} + +TEST_F(AmdsmiWrapMicrotest, LinkInfo_ArsmiFailureIsInternalError) { + InitFor(/*lib=*/false); + g_arsmi.linkResult = 3; + amdsmi_link_type_t type; + int hops = 0, count = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclInternalError); +} + +TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibQueriesTheSrcAndDstGpus) { + InitFor(/*lib=*/true); + TwoGpusBehindACpu(); + g_smi.linkType = AMDSMI_LINK_TYPE_PCIE; + g_smi.linkHops = 3; + amdsmi_link_type_t type; + int hops = 0, count = 0; + ASSERT_EQ(amd_smi_getLinkInfo(1, 0, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(g_smi.lastLinkSrc, &g_smi.sockets[0][2]); + EXPECT_EQ(g_smi.lastLinkDst, &g_smi.sockets[0][1]); + EXPECT_EQ(type, AMDSMI_LINK_TYPE_PCIE); + EXPECT_EQ(hops, 3); + EXPECT_EQ(count, 1); +} + +// XGMI links report weight 0, so their count comes from the max/min bandwidth ratio instead. +TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibXgmiCountIsTheBandwidthRatio) { + InitFor(/*lib=*/true); + TwoGpusBehindACpu(); + g_smi.linkType = AMDSMI_LINK_TYPE_XGMI; + g_smi.minBw = 64; + g_smi.maxBw = 448; + amdsmi_link_type_t type; + int hops = 0, count = 0; + ASSERT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(count, 7); + g_smi.minBw = 0; + ASSERT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(count, 1); +} + +// An unknown index is logged, and the query still runs against the null handle (pinned as observed). +TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibUnknownIndexIsLogged) { + InitFor(/*lib=*/true); + TwoGpusBehindACpu(); + amdsmi_link_type_t type; + int hops = 0, count = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 9, &type, &hops, &count), ncclSuccess); + EXPECT_TRUE(Logged(NCCL_LOG_ERROR, "could not find processor handle for dstIndex: 9")); + EXPECT_FALSE(Logged(NCCL_LOG_ERROR, "srcIndex")); + EXPECT_EQ(g_smi.lastLinkDst, nullptr); +} + +// --------------------------------------------------------------------------- +// amd_smi_getFirmwareVersion +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_InitFailureIsReturned) { + g_arsmi.initResult = 1; + uint64_t fw = 7; + EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclInternalError); + EXPECT_EQ(fw, 7u); +} + +TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_Wsl2IsZero) { + UseWsl2(); + uint64_t fw = 7; + EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclSuccess); + EXPECT_EQ(fw, 0u); +} + +TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_ArsmiReportsItsVersion) { + g_arsmi.fwVersion = 0x55; + uint64_t fw = 0; + EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclSuccess); + EXPECT_EQ(fw, 0x55u); + g_arsmi.fwResult = 2; + EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclInternalError); +} + +TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_LibReportsTheFirstEntry) { + UseLib(); + TwoGpusBehindACpu(); + g_smi.sockets[0][2].fwVersion = 0x99; + uint64_t fw = 0; + EXPECT_EQ(amd_smi_getFirmwareVersion(1, &fw), ncclSuccess); + EXPECT_EQ(fw, 0x99u); +} + +TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_LibWithNoEntriesIsZero) { + UseLib(); + TwoGpusBehindACpu(); + g_smi.sockets[0][1].numFwInfo = 0; + g_smi.sockets[0][1].fwVersion = 0x99; + uint64_t fw = 7; + EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclSuccess); + EXPECT_EQ(fw, 0u); +} + +TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_LibFailures) { + UseLib(); + TwoGpusBehindACpu(); + uint64_t fw = 0; + EXPECT_EQ(amd_smi_getFirmwareVersion(5, &fw), ncclInternalError); // no such GPU + g_smi.fwStatus = AMDSMI_STATUS_INVAL; + EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclInternalError); + g_smi.fwStatus = AMDSMI_STATUS_SUCCESS; + pfn_amdsmi_get_fw_info = nullptr; + EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclInternalError); + EXPECT_TRUE(Logged(NCCL_LOG_ERROR, "amdsmi_get_fw_info symbol not loaded")); +} + +// --------------------------------------------------------------------------- +// amd_smi_ensureFabricInitialized and the fabric accessors +// --------------------------------------------------------------------------- + +namespace { +ARSMI_fabricInfo ArsmiFabric(int supported, uint32_t accelId, uint32_t bw) { + ARSMI_fabricInfo f{}; + f.supported = supported; + f.fabric_type = ARSMI_FABRIC_TYPE_UALLINK; + f.accel_state = ARSMI_FABRIC_ACCELERATOR_VPOD_STATE_ACTIVE; + f.accel_id = accelId; + f.bandwidth = bw; + f.latency = 900; + f.vpod_id = 3; + f.vpod_size = 8; + f.ppod_size = 64; + for (int i = 0; i < 16; ++i) f.ppod_id[i] = static_cast(i + 1); + return f; +} +} // namespace + +TEST_F(AmdsmiWrapMicrotest, Fabric_InitFailureIsCachedForTheThread) { + g_arsmi.initResult = 1; + EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclInternalError); + g_arsmi.initResult = 0; + EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclInternalError); +} + +TEST_F(AmdsmiWrapMicrotest, Fabric_Wsl2HasNoFabric) { + UseWsl2(); + EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_EQ(amdsmiFabricDeviceCount, 0); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "skipped: WSL2")); +} + +TEST_F(AmdsmiWrapMicrotest, Fabric_NoDevicesIsSuccessWithAnEmptyCache) { + EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_EQ(amdsmiFabricDeviceCount, 0); +} + +TEST_F(AmdsmiWrapMicrotest, Fabric_TooManyDevicesIsInternalError) { + g_arsmi.numDevices = amdsmiFabricMaxDevices + 1; + EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclInternalError); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "exceeds internal maximum")); + g_arsmi.numDevices = amdsmiFabricMaxDevices; // the boundary itself fits + Reset(); + g_arsmi.numDevices = amdsmiFabricMaxDevices; + EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_EQ(amdsmiFabricDeviceCount, amdsmiFabricMaxDevices); +} + +TEST_F(AmdsmiWrapMicrotest, Fabric_SysfsPathCopiesEachDevicesFields) { + g_arsmi.numDevices = 2; + g_arsmi.fabric[0] = ArsmiFabric(1, 11, 400000); // device 1 has no ualink dir + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + ASSERT_EQ(amdsmiFabricDeviceCount, 2); + const auto& d0 = amdsmiFabricDevices[0]; + EXPECT_TRUE(d0.fabricSupported); + EXPECT_EQ(d0.fabricType, AMDSMI_FABRIC_TYPE_UALLINK); + EXPECT_EQ(d0.state, AMDSMI_FABRIC_ACCELERATOR_VPOD_STATE_ACTIVE); + EXPECT_EQ(d0.acceleratorId, 11u); + EXPECT_EQ(d0.bandwidth, 400000u); + EXPECT_EQ(d0.latency, 900u); + EXPECT_EQ(d0.cliqueId, 3u); + EXPECT_EQ(d0.vpodSize, 8u); + EXPECT_EQ(d0.ppodSize, 64u); + EXPECT_EQ(d0.clusterUuid[0], 1); + EXPECT_EQ(d0.clusterUuid[15], 16); + EXPECT_FALSE(amdsmiFabricDevices[1].fabricSupported); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "UALLink fabric detected (sysfs)")); +} + +// Discovery runs once; later calls, on any thread, read the cache. +TEST_F(AmdsmiWrapMicrotest, Fabric_DiscoveryRunsOnce) { + g_arsmi.numDevices = 1; + g_arsmi.fabric[0] = ArsmiFabric(1, 11, 100); + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + g_arsmi.fabric[0] = ArsmiFabric(1, 22, 200); + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + std::thread other([] { EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); }); + other.join(); + EXPECT_EQ(amdsmiFabricDevices[0].acceleratorId, 11u); +} + +TEST_F(AmdsmiWrapMicrotest, Fabric_LibPathReadsTheTypedStruct) { + if (!(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) { + GTEST_SKIP() << "this build's amd_smi header uses the extended-union layout, which always reads sysfs"; + } + UseLib(); + TwoGpusBehindACpu(); + g_smi.sockets[0].erase(g_smi.sockets[0].begin()); // GPUs only: getNumDevice counts every processor + FakeProc& g1 = g_smi.sockets[0][1]; + g1.accelId = 21; + g1.bandwidth = 800000; + g1.latency = 700; + g1.vpodId = 5; + g1.vpodSize = 16; + g1.ppodSize = 128; + g1.ppodId[0] = 0xAB; + g1.fabricType = AMDSMI_FABRIC_TYPE_UALLINK; + g_smi.sockets[0][0].accelState = AMDSMI_FABRIC_ACCELERATOR_VPOD_STATE_ERROR; // not usable + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + ASSERT_EQ(amdsmiFabricDeviceCount, 2); + EXPECT_FALSE(amdsmiFabricDevices[0].fabricSupported); + const auto& d1 = amdsmiFabricDevices[1]; + EXPECT_TRUE(d1.fabricSupported); + EXPECT_EQ(d1.fabricType, AMDSMI_FABRIC_TYPE_UALLINK); + EXPECT_EQ(d1.acceleratorId, 21u); + EXPECT_EQ(d1.bandwidth, 800000u); + EXPECT_EQ(d1.latency, 700u); + EXPECT_EQ(d1.cliqueId, 5u); + EXPECT_EQ(d1.vpodSize, 16u); + EXPECT_EQ(d1.ppodSize, 128u); + EXPECT_EQ(d1.clusterUuid[0], 0xAB); + EXPECT_EQ(g_arsmi.fabric.size(), 0u); + EXPECT_FALSE(Logged(NCCL_LOG_WARN, "falling back to sysfs")); + EXPECT_TRUE(Logged(NCCL_LOG_INFO, "GPU 1: UALLink fabric detected - ")); +} + +// A version the library claims but RCCL cannot read is refused; the unreported sentinel reads as v1. +TEST_F(AmdsmiWrapMicrotest, Fabric_LibVersionGate) { + if (!(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) GTEST_SKIP() << "extended-union layout reads sysfs"; + UseLib(); + FakeProc a, b; + a.hipId = 0; + a.fabricVersion = 7; + b.hipId = 1; + b.fabricVersion = kAmdSmiFabricVersionUnreported; + g_smi.sockets = {{a, b}}; + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_FALSE(amdsmiFabricDevices[0].fabricSupported); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "unexpected fabric info version 7 for device 0")); + EXPECT_TRUE(amdsmiFabricDevices[1].fabricSupported); +} + +TEST_F(AmdsmiWrapMicrotest, Fabric_LibPerDeviceFailureMarksOnlyThatDeviceUnsupported) { + if (!(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) GTEST_SKIP() << "extended-union layout reads sysfs"; + UseLib(); + FakeProc a, b; + a.hipId = 0; + b.hipId = 1; + b.fabricOk = false; + g_smi.sockets = {{a, b}}; + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_TRUE(amdsmiFabricDevices[0].fabricSupported); + EXPECT_FALSE(amdsmiFabricDevices[1].fabricSupported); +} + +// Each way the ABI probe can leave the library's layout unconfirmed falls back to sysfs. +TEST_F(AmdsmiWrapMicrotest, Fabric_LibProbeFallsBackToSysfs) { + struct Case { + const char* name; + std::function arrange; + const char* warn; + }; + const size_t otherExtent = amdSmiFabricLayoutIs8Gpu ? kExtent16Gpu : kExtent8Gpu; + const std::vector cases = { + {"probe call fails", [] { g_smi.fabricProbeStatus = AMDSMI_STATUS_NOT_SUPPORTED; }, "the probe call returned"}, + // 300 overruns the 8-GPU object but stops short of the 16-GPU one: no known library writes that. + {"unknown extent", [] { g_smi.fabricExtent = 300; }, "matches no layout RCCL knows"}, + {"extended union", [] { g_smi.fabricExtent = kExtentExtendedUnion; }, "wrote only the v1 payload"}, + {"other layout", [otherExtent] { g_smi.fabricExtent = otherExtent; }, "ABI mismatch"}, + }; + for (const auto& c : cases) { + SCOPED_TRACE(c.name); + if (std::strcmp(c.name, "other layout") == 0 && !(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) continue; + Reset(); + UseLib(); + FakeProc a; + g_smi.sockets = {{a}}; + g_arsmi.fabric[0] = ArsmiFabric(1, 33, 100); + c.arrange(); + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, c.warn)); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "falling back to sysfs")); + EXPECT_EQ(amdsmiFabricDevices[0].acceleratorId, 33u); // came from ARSMI + } +} + +TEST_F(AmdsmiWrapMicrotest, Fabric_LibWithoutAProbeHandleFallsBackToSysfs) { + UseLib(); + FakeProc a; + a.hipId = 4; // no GPU answers to hip id 0 + g_smi.sockets = {{a}}; + g_arsmi.fabric[0] = ArsmiFabric(1, 44, 100); + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "no processor handle to probe")); + EXPECT_EQ(amdsmiFabricDevices[0].acceleratorId, 44u); +} + +// Without the fabric entry point the probe is skipped and each device is marked unsupported. +TEST_F(AmdsmiWrapMicrotest, Fabric_LibWithoutTheFabricSymbolMarksDevicesUnsupported) { + UseLib(); + FakeProc a; + g_smi.sockets = {{a}}; + g_arsmi.fabric[0] = ArsmiFabric(1, 44, 100); + pfn_amdsmi_get_gpu_fabric_info = nullptr; + ASSERT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); + EXPECT_FALSE(amdsmiFabricDevices[0].fabricSupported); + EXPECT_TRUE(Logged(NCCL_LOG_WARN, "fabric functions not loaded for device 0")); +} + +TEST_F(AmdsmiWrapMicrotest, FabricAccessors_ReadTheCacheWithBoundsChecks) { + g_arsmi.numDevices = 2; + g_arsmi.fabric[0] = ArsmiFabric(1, 11, 400000); + g_arsmi.fabric[1] = ArsmiFabric(0, 12, 300000); // present but not usable + + bool supported = false; + ASSERT_EQ(amd_smi_isFabricSupported(0, &supported), ncclSuccess); + EXPECT_TRUE(supported); + ASSERT_EQ(amd_smi_isFabricSupported(1, &supported), ncclSuccess); + EXPECT_FALSE(supported); + supported = true; + ASSERT_EQ(amd_smi_isFabricSupported(2, &supported), ncclSuccess); + EXPECT_FALSE(supported); + + amdsmiFabricDeviceInfo info{}; + ASSERT_EQ(amd_smi_getFabricDeviceInfo(0, &info), ncclSuccess); + EXPECT_EQ(info.acceleratorId, 11u); + EXPECT_EQ(amd_smi_getFabricDeviceInfo(2, &info), ncclInvalidArgument); + + uint32_t bw = 7; + ASSERT_EQ(amd_smi_getFabricBandwidth(0, &bw), ncclSuccess); + EXPECT_EQ(bw, 400000u); + ASSERT_EQ(amd_smi_getFabricBandwidth(1, &bw), ncclSuccess); // unsupported: 0, not its bandwidth + EXPECT_EQ(bw, 0u); + bw = 7; + ASSERT_EQ(amd_smi_getFabricBandwidth(2, &bw), ncclSuccess); + EXPECT_EQ(bw, 0u); +} + +TEST_F(AmdsmiWrapMicrotest, FabricAccessors_PropagateAnInitFailure) { + g_arsmi.initResult = 1; + bool supported = true; + EXPECT_EQ(amd_smi_isFabricSupported(0, &supported), ncclInternalError); + amdsmiFabricDeviceInfo info{}; + EXPECT_EQ(amd_smi_getFabricDeviceInfo(0, &info), ncclInternalError); + uint32_t bw = 0; + EXPECT_EQ(amd_smi_getFabricBandwidth(0, &bw), ncclInternalError); +} + +// --------------------------------------------------------------------------- +// Fabric telemetry +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, Telemetry_IsAmdSmiLibOnly) { + amdsmi_fabric_telemetry_t* t = nullptr; + auto* dummy = reinterpret_cast(0x2000); + EXPECT_EQ(amd_smi_allocFabricTelemetry(0, 1, &t), ncclSystemError); + EXPECT_EQ(amd_smi_getFabricTelemetryData(0, dummy), ncclSystemError); + EXPECT_EQ(amd_smi_freeFabricTelemetry(0, dummy), ncclSystemError); + EXPECT_EQ(g_smi.telemAllocCalls + g_smi.telemGetCalls + g_smi.telemFreeCalls, 0); +} + +TEST_F(AmdsmiWrapMicrotest, Telemetry_NullBufferIsRejected) { + UseLib(); + TwoGpusBehindACpu(); + EXPECT_EQ(amd_smi_getFabricTelemetryData(0, nullptr), ncclSystemError); + EXPECT_EQ(amd_smi_freeFabricTelemetry(0, nullptr), ncclSystemError); + EXPECT_EQ(g_smi.telemGetCalls + g_smi.telemFreeCalls, 0); +} + +TEST_F(AmdsmiWrapMicrotest, Telemetry_CallsTheLibraryWithTheDevicesHandle) { + UseLib(); + TwoGpusBehindACpu(); + amdsmi_fabric_telemetry_t* t = nullptr; + ASSERT_EQ(amd_smi_allocFabricTelemetry(1, 0x6, &t), ncclSuccess); + EXPECT_EQ(t, reinterpret_cast(0x1000)); + EXPECT_EQ(g_smi.lastTelemHandle, &g_smi.sockets[0][2]); + EXPECT_EQ(g_smi.lastTelemMask, 0x6u); + g_smi.lastTelemHandle = nullptr; + ASSERT_EQ(amd_smi_getFabricTelemetryData(1, t), ncclSuccess); + EXPECT_EQ(g_smi.lastTelemHandle, &g_smi.sockets[0][2]); + g_smi.lastTelemHandle = nullptr; + ASSERT_EQ(amd_smi_freeFabricTelemetry(1, t), ncclSuccess); + EXPECT_EQ(g_smi.lastTelemHandle, &g_smi.sockets[0][2]); +} + +TEST_F(AmdsmiWrapMicrotest, Telemetry_FailuresAreReturned) { + UseLib(); + TwoGpusBehindACpu(); + amdsmi_fabric_telemetry_t* t = nullptr; + auto* dummy = reinterpret_cast(0x2000); + EXPECT_EQ(amd_smi_allocFabricTelemetry(9, 1, &t), ncclInternalError); // no such GPU + EXPECT_EQ(amd_smi_getFabricTelemetryData(9, dummy), ncclInternalError); + EXPECT_EQ(amd_smi_freeFabricTelemetry(9, dummy), ncclInternalError); + g_smi.telemStatus = AMDSMI_STATUS_INVAL; + EXPECT_EQ(amd_smi_allocFabricTelemetry(0, 1, &t), ncclInternalError); + EXPECT_EQ(amd_smi_getFabricTelemetryData(0, dummy), ncclInternalError); + EXPECT_EQ(amd_smi_freeFabricTelemetry(0, dummy), ncclInternalError); +} + +// --------------------------------------------------------------------------- +// amd_smi_fabricTelemIdToString +// --------------------------------------------------------------------------- + +TEST_F(AmdsmiWrapMicrotest, TelemIdToString_WithoutTheSymbolAsksArsmi) { + pfn_amdsmi_fabric_telem_id_to_string = nullptr; + EXPECT_STREQ(amd_smi_fabricTelemIdToString(5), "ARSMI_NAME"); +} + +// amd_smi 27+ and an unknown version use the out-parameter form. +TEST_F(AmdsmiWrapMicrotest, TelemIdToString_OutParamForm) { + for (uint32_t major : {kAmdSmiLibVersionUnknown, 27u, 30u}) { + SCOPED_TRACE(major); + amdSmiLibMajor.store(major); + g_smi.telemNameStatus = AMDSMI_STATUS_SUCCESS; + g_smi.telemName = "OUTPARAM_NAME"; + EXPECT_STREQ(amd_smi_fabricTelemIdToString(5), "OUTPARAM_NAME"); + g_smi.telemNameStatus = AMDSMI_STATUS_INVAL; + EXPECT_STREQ(amd_smi_fabricTelemIdToString(5), "UNKNOWN"); + g_smi.telemNameStatus = AMDSMI_STATUS_SUCCESS; + g_smi.telemName = nullptr; + EXPECT_STREQ(amd_smi_fabricTelemIdToString(5), "UNKNOWN"); + } +} + +// Before 27 the symbol returns the name, so it is called through that signature instead. +TEST_F(AmdsmiWrapMicrotest, TelemIdToString_LegacyFormBeforeMajor27) { + amdSmiLibMajor.store(26); + pfn_amdsmi_fabric_telem_id_to_string = + reinterpret_cast(FakeLegacyTelemIdToString); + g_legacyTelemName = "LEGACY_NAME"; + EXPECT_STREQ(amd_smi_fabricTelemIdToString(5), "LEGACY_NAME"); + g_legacyTelemName = nullptr; + EXPECT_STREQ(amd_smi_fabricTelemIdToString(5), "UNKNOWN"); +} diff --git a/projects/rccl/test/host/run_host_tests.sh b/projects/rccl/test/host/run_host_tests.sh index ec702f4f62b1..10cb49448e0d 100755 --- a/projects/rccl/test/host/run_host_tests.sh +++ b/projects/rccl/test/host/run_host_tests.sh @@ -168,6 +168,7 @@ do_host_tests() { "rccl-UnitTestsMicroCe:$SCRIPT_DIR/host_tests_micro_ce.xml" "rccl-UnitTestsMicroClientSupport:$SCRIPT_DIR/host_tests_micro_client_support.xml" "rccl-UnitTestsMicroGinHost:$SCRIPT_DIR/host_tests_micro_gin_host.xml" + "rccl-UnitTestsMicroAmdsmiWrap:$SCRIPT_DIR/host_tests_micro_amdsmi_wrap.xml" "rccl-UnitTestsMicroDiagnostics:$SCRIPT_DIR/host_tests_micro_diagnostics.xml" ) # Binaries that only exist for some CMake option settings (rccl-UnitTestsMicroSymKernels needs diff --git a/projects/rccl/test/test_categories_micro_amdsmi_wrap.yaml b/projects/rccl/test/test_categories_micro_amdsmi_wrap.yaml new file mode 100644 index 000000000000..7fc99ae362f7 --- /dev/null +++ b/projects/rccl/test/test_categories_micro_amdsmi_wrap.yaml @@ -0,0 +1,47 @@ +# rccl-UnitTestsMicroAmdsmiWrap — shared/ctest schema. +# Host-only micro-tests for src/misc/amdsmi_wrap.cc -- GPU enumeration, bus IDs, link topology, +# firmware and UALoE/UALink fabric discovery over all three backends (no GPU, no librccl). +# RCCL does not support Windows: exclude_windows matches all tests so generated suites are empty on Windows. + +test_categories: + quick: µ_amdsmi_wrap_quick + description: "Smoke — host-only src/misc/amdsmi_wrap.cc micro-tests, sub-second." + test_patterns: + - "AmdsmiWrap*" + exclude_windows: + - "*" + labels: + - "rccl" + - "quick" + - "smoke" + - "precheck" + - "microtest" + standard: + <<: *micro_amdsmi_wrap_quick + labels: + - "rccl" + - "standard" + - "regression" + - "microtest" + comprehensive: + <<: *micro_amdsmi_wrap_quick + labels: + - "rccl" + - "comprehensive" + - "comprehensive_regression" + - "microtest" + full: + <<: *micro_amdsmi_wrap_quick + labels: + - "rccl" + - "full" + - "full_regression" + - "microtest" + +execution_settings: + timeout_multiplier: 1 + category_timeouts: + quick: 600 + standard: 600 + comprehensive: 600 + full: 600 diff --git a/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json b/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json index 07da6b28ba73..962ee52b1a33 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json +++ b/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json @@ -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_amdsmi_wrap": { + "is_gtest": true, + "binary": "host/rccl-UnitTestsMicroAmdsmiWrap", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 60, + "tests": [ + {"name": "AmdsmiWrapMicrotests", "description": "Host-only misc/amdsmi_wrap.cc microtests (no GPU, no HIP/RCCL runtime); covers device, bus-id, link, firmware and fabric queries across the WSL2, amd_smi_lib and ARSMI backends", "test_filter": "*"} + ] } }, "test_suites": [ @@ -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 - amdsmi_wrap.cc", + "description": "rccl-UnitTestsMicroAmdsmiWrap (host-only misc/amdsmi_wrap.cc microtests, no GPU/HIP/RCCL runtime)", + "config": "unit_tests_micro_amdsmi_wrap", + "enabled": true } ] } diff --git a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py index 70add7053ab2..00df046bb25e 100644 --- a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py +++ b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py @@ -3292,7 +3292,8 @@ def generate_coverage_report(self): "rccl-UnitTestsMicroInit-faultinj", "rccl-UnitTestsMicroEnqueue", "rccl-UnitTestsMicroSymKernels", - "rccl-UnitTestsMicroGinHost"]: + "rccl-UnitTestsMicroGinHost", + "rccl-UnitTestsMicroAmdsmiWrap"]: binary_path = os.path.join(host_test_dir, binary) if os.path.isfile(binary_path): object_files.extend(["--object", binary_path]) From 143d2eaff62a8f894564d721dac5efef7a21e045 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 14:04:43 -0700 Subject: [PATCH 09/16] test(rccl): tighten amdsmi_wrap microtests per review Seed a nonzero firmware version in the WSL2 test so only the early return can produce 0, drive install/clear/dlsym from one symbol table and check every pfn slot after dlopen, count calls to prove getProcessorHandle's early returns, name the typed-fabric-layout condition once, and drop a dead store in the too-many-devices test. Co-Authored-By: Claude --- projects/rccl/test/host/amdsmi-wrap-test.cc | 137 ++++++++++---------- 1 file changed, 67 insertions(+), 70 deletions(-) diff --git a/projects/rccl/test/host/amdsmi-wrap-test.cc b/projects/rccl/test/host/amdsmi-wrap-test.cc index c176b8aba7b9..0f92d1c82faa 100644 --- a/projects/rccl/test/host/amdsmi-wrap-test.cc +++ b/projects/rccl/test/host/amdsmi-wrap-test.cc @@ -172,6 +172,10 @@ constexpr size_t kExtentExtendedUnion = kAmdSmiFabricV1PayloadEnd; constexpr size_t kCompiledExtent = amdSmiFabricLayoutIs8Gpu ? kExtent8Gpu : (amdSmiFabricLayoutIs16Gpu ? kExtent16Gpu : kExtentExtendedUnion); +// The typed amd_smi fabric path is only ever taken for the 8- and 16-GPU layouts; the extended +// union always falls back to sysfs, so tests of the typed path skip on such a build. +constexpr bool kTypedFabricLayout = amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu; + template void SetFabricVersion(FabricInfoT& info, uint32_t version) { if constexpr (amdSmiFabricInfoIsFlat::value) { @@ -224,6 +228,9 @@ struct FakeSmi { uint64_t maxBw = 0; int initCalls = 0; int shutdownCalls = 0; + int procsCalls = 0; + int typeCalls = 0; + int enumCalls = 0; int fabricInfoCalls = 0; int telemAllocCalls = 0; int telemGetCalls = 0; @@ -269,6 +276,7 @@ amdsmi_status_t FakeSocketHandles(uint32_t* count, amdsmi_socket_handle* handles return AMDSMI_STATUS_SUCCESS; } amdsmi_status_t FakeProcessorHandles(amdsmi_socket_handle socket, uint32_t* count, amdsmi_processor_handle* handles) { + ++g_smi.procsCalls; if (g_smi.procsStatus != AMDSMI_STATUS_SUCCESS) return g_smi.procsStatus; auto& procs = *static_cast*>(socket); if (handles != nullptr) { @@ -278,10 +286,12 @@ amdsmi_status_t FakeProcessorHandles(amdsmi_socket_handle socket, uint32_t* coun return AMDSMI_STATUS_SUCCESS; } amdsmi_status_t FakeProcessorType(amdsmi_processor_handle h, processor_type_t* type) { + ++g_smi.typeCalls; *type = Proc(h)->gpu ? AMDSMI_PROCESSOR_TYPE_AMD_GPU : AMDSMI_PROCESSOR_TYPE_AMD_CPU; return AMDSMI_STATUS_SUCCESS; } amdsmi_status_t FakeEnumerationInfo(amdsmi_processor_handle h, amdsmi_enumeration_info_t* info) { + ++g_smi.enumCalls; if (g_smi.enumStatus != AMDSMI_STATUS_SUCCESS) return g_smi.enumStatus; std::memset(info, 0, sizeof(*info)); info->hip_id = Proc(h)->hipId; @@ -378,48 +388,46 @@ amdsmi_status_t FakeTelemIdToString(uint64_t, const char** name) { const char* g_legacyTelemName = "LEGACY_NAME"; const char* FakeLegacyTelemIdToString(uint64_t) { return g_legacyTelemName; } +// Every amd_smi entry point amdsmi_wrap.cc resolves, by the name it dlsym()s, with its pfn_* slot and +// this file's fake. One list drives installing the fakes, clearing them, and the dlopen test's dlsym. +struct SmiSymbol { + const char* name; + void** slot; + void* fake; +}; +#define SMI_SYMBOL(fn, fake) SmiSymbol{#fn, reinterpret_cast(&pfn_##fn), reinterpret_cast(fake)} +const std::vector& SmiSymbols() { + static const std::vector symbols = { + SMI_SYMBOL(amdsmi_init, FakeInit), + SMI_SYMBOL(amdsmi_shut_down, FakeShutDown), + SMI_SYMBOL(amdsmi_status_code_to_string, FakeStatusString), + SMI_SYMBOL(amdsmi_get_lib_version, FakeLibVersion), + SMI_SYMBOL(amdsmi_get_socket_handles, FakeSocketHandles), + SMI_SYMBOL(amdsmi_get_processor_handles, FakeProcessorHandles), + SMI_SYMBOL(amdsmi_get_processor_type, FakeProcessorType), + SMI_SYMBOL(amdsmi_get_processor_handle_from_bdf, FakeHandleFromBdf), + SMI_SYMBOL(amdsmi_get_gpu_enumeration_info, FakeEnumerationInfo), + SMI_SYMBOL(amdsmi_get_gpu_bdf_id, FakeBdfId), + SMI_SYMBOL(amdsmi_topo_get_link_type, FakeLinkType), + SMI_SYMBOL(amdsmi_topo_get_link_weight, FakeLinkWeight), + SMI_SYMBOL(amdsmi_get_minmax_bandwidth_between_processors, FakeMinMaxBw), + SMI_SYMBOL(amdsmi_get_gpu_fabric_info, FakeFabricInfo), + SMI_SYMBOL(amdsmi_alloc_fabric_telemetry, FakeAllocTelem), + SMI_SYMBOL(amdsmi_get_fabric_telemetry_data, FakeGetTelem), + SMI_SYMBOL(amdsmi_free_fabric_telemetry, FakeFreeTelem), + SMI_SYMBOL(amdsmi_fabric_telem_id_to_string, FakeTelemIdToString), + SMI_SYMBOL(amdsmi_get_fw_info, FakeFwInfo), + }; + return symbols; +} +#undef SMI_SYMBOL + void InstallFakeSmiTable() { - pfn_amdsmi_init = FakeInit; - pfn_amdsmi_shut_down = FakeShutDown; - pfn_amdsmi_status_code_to_string = FakeStatusString; - pfn_amdsmi_get_lib_version = FakeLibVersion; - pfn_amdsmi_get_socket_handles = FakeSocketHandles; - pfn_amdsmi_get_processor_handles = FakeProcessorHandles; - pfn_amdsmi_get_processor_type = FakeProcessorType; - pfn_amdsmi_get_processor_handle_from_bdf = FakeHandleFromBdf; - pfn_amdsmi_get_gpu_enumeration_info = FakeEnumerationInfo; - pfn_amdsmi_get_gpu_bdf_id = FakeBdfId; - pfn_amdsmi_topo_get_link_type = FakeLinkType; - pfn_amdsmi_topo_get_link_weight = FakeLinkWeight; - pfn_amdsmi_get_minmax_bandwidth_between_processors = FakeMinMaxBw; - pfn_amdsmi_get_gpu_fabric_info = FakeFabricInfo; - pfn_amdsmi_alloc_fabric_telemetry = FakeAllocTelem; - pfn_amdsmi_get_fabric_telemetry_data = FakeGetTelem; - pfn_amdsmi_free_fabric_telemetry = FakeFreeTelem; - pfn_amdsmi_fabric_telem_id_to_string = FakeTelemIdToString; - pfn_amdsmi_get_fw_info = FakeFwInfo; + for (const auto& sym : SmiSymbols()) *sym.slot = sym.fake; } void ClearSmiTable() { - pfn_amdsmi_init = nullptr; - pfn_amdsmi_shut_down = nullptr; - pfn_amdsmi_status_code_to_string = nullptr; - pfn_amdsmi_get_lib_version = nullptr; - pfn_amdsmi_get_socket_handles = nullptr; - pfn_amdsmi_get_processor_handles = nullptr; - pfn_amdsmi_get_processor_type = nullptr; - pfn_amdsmi_get_processor_handle_from_bdf = nullptr; - pfn_amdsmi_get_gpu_enumeration_info = nullptr; - pfn_amdsmi_get_gpu_bdf_id = nullptr; - pfn_amdsmi_topo_get_link_type = nullptr; - pfn_amdsmi_topo_get_link_weight = nullptr; - pfn_amdsmi_get_minmax_bandwidth_between_processors = nullptr; - pfn_amdsmi_get_gpu_fabric_info = nullptr; - pfn_amdsmi_alloc_fabric_telemetry = nullptr; - pfn_amdsmi_get_fabric_telemetry_data = nullptr; - pfn_amdsmi_free_fabric_telemetry = nullptr; - pfn_amdsmi_fabric_telem_id_to_string = nullptr; - pfn_amdsmi_get_fw_info = nullptr; + for (const auto& sym : SmiSymbols()) *sym.slot = nullptr; } class AmdsmiWrapMicrotest : public ::testing::Test { @@ -615,32 +623,16 @@ int DlopenResolvesChild() { g_params["RCCL_USE_AMD_SMI_LIB"] = 1; static int handle; g_dlopen = [](const char*, int) -> void* { return &handle; }; - g_dlsymTable = { - {"amdsmi_init", reinterpret_cast(FakeInit)}, - {"amdsmi_shut_down", reinterpret_cast(FakeShutDown)}, - {"amdsmi_status_code_to_string", reinterpret_cast(FakeStatusString)}, - {"amdsmi_get_lib_version", reinterpret_cast(FakeLibVersion)}, - {"amdsmi_get_socket_handles", reinterpret_cast(FakeSocketHandles)}, - {"amdsmi_get_processor_handles", reinterpret_cast(FakeProcessorHandles)}, - {"amdsmi_get_processor_type", reinterpret_cast(FakeProcessorType)}, - {"amdsmi_get_processor_handle_from_bdf", reinterpret_cast(FakeHandleFromBdf)}, - {"amdsmi_get_gpu_enumeration_info", reinterpret_cast(FakeEnumerationInfo)}, - {"amdsmi_get_gpu_bdf_id", reinterpret_cast(FakeBdfId)}, - {"amdsmi_topo_get_link_type", reinterpret_cast(FakeLinkType)}, - {"amdsmi_topo_get_link_weight", reinterpret_cast(FakeLinkWeight)}, - {"amdsmi_get_minmax_bandwidth_between_processors", reinterpret_cast(FakeMinMaxBw)}, - {"amdsmi_get_gpu_fabric_info", reinterpret_cast(FakeFabricInfo)}, - {"amdsmi_fabric_telem_id_to_string", reinterpret_cast(FakeTelemIdToString)}, - {"amdsmi_alloc_fabric_telemetry", reinterpret_cast(FakeAllocTelem)}, - {"amdsmi_get_fabric_telemetry_data", reinterpret_cast(FakeGetTelem)}, - {"amdsmi_free_fabric_telemetry", reinterpret_cast(FakeFreeTelem)}, - {"amdsmi_get_fw_info", reinterpret_cast(FakeFwInfo)}, - }; - bool ok = amd_smi_init() == ncclSuccess && g_smi.initCalls == 1; - ok = ok && pfn_amdsmi_get_gpu_fabric_info == FakeFabricInfo && pfn_amdsmi_get_fw_info == FakeFwInfo && - pfn_amdsmi_get_minmax_bandwidth_between_processors == FakeMinMaxBw && - pfn_amdsmi_free_fabric_telemetry == FakeFreeTelem && pfn_amdsmi_get_lib_version == FakeLibVersion; - return ok ? 0 : 1; + for (const auto& sym : SmiSymbols()) g_dlsymTable[sym.name] = sym.fake; + if (amd_smi_init() != ncclSuccess || g_smi.initCalls != 1) return 1; + // Every slot must hold the fake dlsym returned for its own name: a misspelled name leaves it null. + for (const auto& sym : SmiSymbols()) { + if (*sym.slot != sym.fake) { + std::fprintf(stderr, "pfn slot for %s not resolved\n", sym.name); + return 2; + } + } + return 0; } } // namespace @@ -685,11 +677,16 @@ TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_LibraryFailuresAreInternalError) { UseLib(); TwoGpusBehindACpu(); amdsmi_processor_handle h = nullptr; + // Each failure returns at once: nothing past the failing call runs. g_smi.socketsStatus = AMDSMI_STATUS_INVAL; EXPECT_EQ(getProcessorHandle(0, &h), ncclInternalError); + EXPECT_EQ(g_smi.procsCalls, 0); g_smi.socketsStatus = AMDSMI_STATUS_SUCCESS; g_smi.enumStatus = AMDSMI_STATUS_INVAL; EXPECT_EQ(getProcessorHandle(0, &h), ncclInternalError); + EXPECT_EQ(g_smi.enumCalls, 1); // the first GPU's query failed; the second GPU was never asked + EXPECT_EQ(g_smi.typeCalls, 2); // the CPU and that first GPU + EXPECT_EQ(h, nullptr); } #if !AMDSMI_DIRECT @@ -1034,6 +1031,7 @@ TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_InitFailureIsReturned) { TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_Wsl2IsZero) { UseWsl2(); + g_arsmi.fwVersion = 0x55; // so only the WSL2 early return can produce 0 uint64_t fw = 7; EXPECT_EQ(amd_smi_getFirmwareVersion(0, &fw), ncclSuccess); EXPECT_EQ(fw, 0u); @@ -1124,8 +1122,7 @@ TEST_F(AmdsmiWrapMicrotest, Fabric_TooManyDevicesIsInternalError) { g_arsmi.numDevices = amdsmiFabricMaxDevices + 1; EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclInternalError); EXPECT_TRUE(Logged(NCCL_LOG_WARN, "exceeds internal maximum")); - g_arsmi.numDevices = amdsmiFabricMaxDevices; // the boundary itself fits - Reset(); + Reset(); // the boundary itself fits g_arsmi.numDevices = amdsmiFabricMaxDevices; EXPECT_EQ(amd_smi_ensureFabricInitialized(), ncclSuccess); EXPECT_EQ(amdsmiFabricDeviceCount, amdsmiFabricMaxDevices); @@ -1165,7 +1162,7 @@ TEST_F(AmdsmiWrapMicrotest, Fabric_DiscoveryRunsOnce) { } TEST_F(AmdsmiWrapMicrotest, Fabric_LibPathReadsTheTypedStruct) { - if (!(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) { + if (!kTypedFabricLayout) { GTEST_SKIP() << "this build's amd_smi header uses the extended-union layout, which always reads sysfs"; } UseLib(); @@ -1201,7 +1198,7 @@ TEST_F(AmdsmiWrapMicrotest, Fabric_LibPathReadsTheTypedStruct) { // A version the library claims but RCCL cannot read is refused; the unreported sentinel reads as v1. TEST_F(AmdsmiWrapMicrotest, Fabric_LibVersionGate) { - if (!(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) GTEST_SKIP() << "extended-union layout reads sysfs"; + if (!kTypedFabricLayout) GTEST_SKIP() << "extended-union layout reads sysfs"; UseLib(); FakeProc a, b; a.hipId = 0; @@ -1216,7 +1213,7 @@ TEST_F(AmdsmiWrapMicrotest, Fabric_LibVersionGate) { } TEST_F(AmdsmiWrapMicrotest, Fabric_LibPerDeviceFailureMarksOnlyThatDeviceUnsupported) { - if (!(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) GTEST_SKIP() << "extended-union layout reads sysfs"; + if (!kTypedFabricLayout) GTEST_SKIP() << "extended-union layout reads sysfs"; UseLib(); FakeProc a, b; a.hipId = 0; @@ -1245,7 +1242,7 @@ TEST_F(AmdsmiWrapMicrotest, Fabric_LibProbeFallsBackToSysfs) { }; for (const auto& c : cases) { SCOPED_TRACE(c.name); - if (std::strcmp(c.name, "other layout") == 0 && !(amdSmiFabricLayoutIs8Gpu || amdSmiFabricLayoutIs16Gpu)) continue; + if (std::strcmp(c.name, "other layout") == 0 && !kTypedFabricLayout) continue; Reset(); UseLib(); FakeProc a; From 05c89f6d5d2fae25600e65f6464b742229e20915 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 14:56:46 -0700 Subject: [PATCH 10/16] test(rccl): pin the XGMI-only gates in amdsmi_wrap link info Give the lib PCIe case a nonzero bandwidth pair and add an ARSMI PCIe link with a direct-link weight, so both XGMI type checks are needed. Co-Authored-By: Claude --- projects/rccl/test/host/amdsmi-wrap-test.cc | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/projects/rccl/test/host/amdsmi-wrap-test.cc b/projects/rccl/test/host/amdsmi-wrap-test.cc index 0f92d1c82faa..44392722462d 100644 --- a/projects/rccl/test/host/amdsmi-wrap-test.cc +++ b/projects/rccl/test/host/amdsmi-wrap-test.cc @@ -965,6 +965,15 @@ TEST_F(AmdsmiWrapMicrotest, LinkInfo_ArsmiDirectXgmiWeightsAreOneHop) { EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); EXPECT_EQ(hops, 1); EXPECT_EQ(count, 1); + // A direct-link weight only counts on XGMI: a PCIe link with weight 15 stays two hops, count 1. + g_arsmi.link.type = ARSMI_IOLINK_TYPE_PCIEXPRESS; + g_arsmi.link.min_bandwidth = 50; + hops = 0; + count = 0; + EXPECT_EQ(amd_smi_getLinkInfo(0, 1, &type, &hops, &count), ncclSuccess); + EXPECT_EQ(type, AMDSMI_LINK_TYPE_PCIE); + EXPECT_EQ(hops, 2); + EXPECT_EQ(count, 1); } TEST_F(AmdsmiWrapMicrotest, LinkInfo_ArsmiFailureIsInternalError) { @@ -980,6 +989,8 @@ TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibQueriesTheSrcAndDstGpus) { TwoGpusBehindACpu(); g_smi.linkType = AMDSMI_LINK_TYPE_PCIE; g_smi.linkHops = 3; + g_smi.minBw = 64; // only XGMI uses the bandwidth ratio, so a PCIe link's count stays 1 + g_smi.maxBw = 448; amdsmi_link_type_t type; int hops = 0, count = 0; ASSERT_EQ(amd_smi_getLinkInfo(1, 0, &type, &hops, &count), ncclSuccess); From cd11a892ec76f46c3cf0aa0e9f90ab7cf13e834f Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 15:40:30 -0700 Subject: [PATCH 11/16] test(rccl): trim microtest docs and comments to match the other host tests Shorten the MICROTEST_README entry, CMake comments and test-file comments to the length the other host microtests use. Comment and doc changes only. Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 14 +--- projects/rccl/test/host/MICROTEST_README.md | 24 ++---- .../rccl/test/host/roctx-nvml-stub-test.cc | 84 ++++--------------- 3 files changed, 28 insertions(+), 94 deletions(-) diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index bd779d79e25b..82f11ad412aa 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -293,11 +293,8 @@ function(rccl_define_micro_source_lists) ../common/main_altrsmi.cpp PARENT_SCOPE) - # rccl-UnitTestsMicroRoctxNvml -- src/misc/roctx.cc + src/misc/nvmlwrap_stub.cc. Its own binary: - # fakes/comm_fakes.cc (linked by rccl-UnitTestsMicro and most other targets) stubs the - # roctx_scoped_range_in constructor/destructor this target compiles as the real thing, and - # ras-diagnostics-gpu-test.cc defines its own NVML seams. The test TU supplies ncclDebugLog and - # g_loadParam itself, so no shared fakes are linked. + # rccl-UnitTestsMicroRoctxNvml -- src/misc/roctx.cc + src/misc/nvmlwrap_stub.cc. Own binary: + # fakes/comm_fakes.cc stubs the roctx_scoped_range_in this target compiles for real. set(TEST_MICRO_ROCTX_NVML_SOURCE_FILES roctx-nvml-stub-test.cc ../common/main_altrsmi.cpp @@ -753,8 +750,7 @@ function(rccl_standalone_setup) # binary; run_host_tests.sh (the only in-repo workflow that runs these) uses # this standalone path, so build the shipping arm explicitly too. set(_ENQ_BUILD_DEVLINKER TRUE PARENT_SCOPE) - # Same reason for roctx.cc's ROCTX_ENABLE arm: in-tree inherits it from the rccl target when - # roctracer is found, so standalone builds that arm explicitly. + # Likewise roctx.cc's ROCTX_ENABLE arm (-roctx). set(_MICRO_BUILD_ROCTX_VARIANT TRUE PARENT_SCOPE) # No GENERATE_SYM_KERNELS variable exists in this standalone project(); check # for the generated host TU on disk instead. @@ -1081,9 +1077,7 @@ set(_roctx_nvml_common_args CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml") rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml ${_roctx_nvml_common_args}) -# roctx.cc gates its roctracer calls on ROCTX_ENABLE at the preprocessor, so one compile covers -# one arm. The -roctx variant covers the roctracer arm against a stand-in -# (fakes/roctx_include), so it builds whether or not roctracer is installed. +# -roctx builds roctx.cc's ROCTX_ENABLE arm against the stand-in fakes/roctx_include. if(_MICRO_BUILD_ROCTX_VARIANT) rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml-roctx ${_roctx_nvml_common_args}) target_compile_definitions(rccl-UnitTestsMicroRoctxNvml-roctx PRIVATE ROCTX_ENABLE) diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index 8a3573f40f06..295e8b70ffa3 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -249,20 +249,9 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: the production source include, after system headers and before test helpers. CMake pins the copied `ncclFuncStr` table to `src/init.cc`. - **`rccl-UnitTestsMicroRoctxNvml`** (+ **`-roctx`**) — `src/misc/roctx.cc` and - `src/misc/nvmlwrap_stub.cc` (via `ROCTX_CC_PATH`/`NVMLWRAP_STUB_CC_PATH`, both - from `roctx-nvml-stub-test.cc`); suites `RoctxNvmlRoctxMicrotest.*` and - `RoctxNvmlStubMicrotest.*`, both caught by the `RoctxNvml*` pattern. Its own - binary because `fakes/comm_fakes.cc`, linked by most other targets, stubs the - `roctx_scoped_range_in` constructor/destructor this file compiles for real, - and `ras-diagnostics-gpu-test.cc` defines its own NVML seams; the TU supplies - its own `ncclDebugLog` and `g_loadParam`, so no shared fakes are linked. - `roctx.cc` gates its `roctxRangePushA`/`roctxRangePop` calls on - `ROCTX_ENABLE` at the preprocessor. Standalone, the base binary covers the - WARN fallback and the standalone-only `-roctx` variant defines `ROCTX_ENABLE` - against the stand-in `fakes/roctx_include/roctracer/roctx.h`, with the test TU - recording each push and pop. In-tree, the base binary inherits `ROCTX_ENABLE` - from the `rccl` target whenever roctracer is found, so it compiles the - roctracer arm instead and no in-tree binary builds the WARN fallback. + `src/misc/nvmlwrap_stub.cc` (via `ROCTX_CC_PATH` / `NVMLWRAP_STUB_CC_PATH`, + from `roctx-nvml-stub-test.cc`); suites `RoctxNvml*`. Its own binary because + `fakes/comm_fakes.cc` stubs the `roctx_scoped_range_in` this file defines. - **`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.*`, @@ -852,13 +841,12 @@ sources. Four of those names are preprocessor variants, not duplicates. `init.cc` gates part of its allocation path on `HIP_*_UNCACHED_MEMORY` and its fault-mask blocks on `ENABLE_FAULT_INJECTION`, and `enqueue.cc` gates `rcclShmemDynamicSize` on -`RCCL_DEVICE_LINKER`, and `roctx.cc` gates its roctracer calls on `ROCTX_ENABLE`, -all at the **preprocessor**, so one compile can only ever reach one arm. The in-RCCL-build path inherits `RCCL_DEVICE_LINKER` from the +`RCCL_DEVICE_LINKER`, and `roctx.cc` gates roctracer on `ROCTX_ENABLE`, all at +the **preprocessor**, so one compile can only ever reach one arm. The in-RCCL-build path inherits `RCCL_DEVICE_LINKER` from the `rccl` target's compile definitions (`ENABLE_DEVICE_LINKER` defaults ON, so the device-linker arm is the one that ships); this standalone project has no `rccl` target to inherit from, which is why it builds the `-devlinker` variant -explicitly, and likewise the `-roctx` variant for `ROCTX_ENABLE`, which in-tree -is inherited only when roctracer is found. `ENABLE_FAULT_INJECTION` was the same inheritance in reverse - +explicitly, and likewise `-roctx`. `ENABLE_FAULT_INJECTION` was the same inheritance in reverse - defined in-tree, absent standalone, so the two paths tested different code - and is now stripped from the inherited list and set per variant in both paths, so the `-faultinj` binary is the only one that has it either way. diff --git a/projects/rccl/test/host/roctx-nvml-stub-test.cc b/projects/rccl/test/host/roctx-nvml-stub-test.cc index c5a295c72315..9835b0756f74 100644 --- a/projects/rccl/test/host/roctx-nvml-stub-test.cc +++ b/projects/rccl/test/host/roctx-nvml-stub-test.cc @@ -4,27 +4,9 @@ * See LICENSE.txt for license information ************************************************************************/ -// Host-only microtests for src/misc/roctx.cc (NVTX payload -> roctx message -// translation and the RAII range used by NCCL_NVTX3_FUNC_RANGE*) and -// src/misc/nvmlwrap_stub.cc (the ROCm stand-in for NVML, which reports no -// devices so NVML-driven callers skip their checks). -// -// Both files are compiled for real into this binary. Neither needs the shared -// fakes: this TU supplies its own ncclDebugLog (to record WARNs) and -// g_loadParam (param_redirect.h routes rcclParamLogRoctx() through it on every -// call instead of caching, so each test picks LOG_ROCTX independently). -// -// The scoped range has two production arms. With ROCTX_ENABLE it calls -// roctxRangePushA/Pop, which this TU defines to record the messages pushed; -// without it, it falls back to one WARN per enter/exit. In-tree builds inherit -// ROCTX_ENABLE from the rccl target when roctracer is found. The standalone -// build compiles this TU twice: rccl-UnitTestsMicroRoctxNvml without it, and -// rccl-UnitTestsMicroRoctxNvml-roctx with it against the stand-in -// fakes/roctracer/roctx.h, so CI runs both arms. The range tests hold in either. -// -// Every stringify() case here stays under MAX_MESSAGE_LENGTH on purpose: past -// it, stringify() writes beyond its 1024-byte buffer (confirmed under ASan), -// tracked as AICOMRCCL-2914 rather than pinned here as undefined behaviour. +// Host-only microtests for src/misc/roctx.cc and src/misc/nvmlwrap_stub.cc. The -roctx variant +// builds roctx.cc's ROCTX_ENABLE arm; the range tests hold in either. stringify() inputs stay +// under MAX_MESSAGE_LENGTH: past it, it overflows its buffer (AICOMRCCL-2914). #include @@ -61,18 +43,14 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long /*flags*/, const char* #include "fakes/param_redirect.h" -// malloc seam for roctx.cc only, so roctxAlloc's two failure paths are reachable. Every header -// roctx.cc includes is pulled in first, so their include guards keep the macro below from -// rewriting any declaration; it applies to roctx.cc's own function bodies and nothing else. +// malloc seam for roctx.cc's own bodies only, so roctxAlloc's failure paths are reachable. #include #include "roctx.h" namespace { std::function g_roctxMalloc; } // namespace -// With ROCTX_ENABLE (inherited in-tree when roctracer is found, or set by the standalone -roctx -// variant), roctx.cc calls the roctracer entry points. Nothing on the micro-test link line provides -// them, so record the calls instead; that also lets the range tests check the exact message pushed. +// With ROCTX_ENABLE, record roctracer's push/pop instead of linking it. #ifdef ROCTX_ENABLE namespace { std::vector g_roctxPushed; @@ -105,8 +83,7 @@ class RoctxNvmlRoctxMicrotest : public ::testing::Test { void TearDown() override { g_loadParam = [](const char*, int64_t deftVal) { return deftVal; }; g_roctxMalloc = nullptr; - // nvtxToRoctx is a production global that two tests insert into; restoring it here rather - // than inline keeps a failed assertion in one of them from leaking into the next test. + // Restored here so a failed assertion can't leak entries into the next test. nvtxToRoctx = savedNvtxToRoctx_; ResetRecorders(); } @@ -119,8 +96,7 @@ class RoctxNvmlRoctxMicrotest : public ::testing::Test { #endif } - // Range enter/exit events, whichever form this build gives them: roctracer push/pop calls with - // ROCTX_ENABLE, one WARN each without it. With logging off there must be none either way. + // Enter/exit events: roctracer push/pop with ROCTX_ENABLE, one WARN each without; none if off. static size_t RangeEvents() { #ifdef ROCTX_ENABLE EXPECT_TRUE(g_logged.empty()) << "a ROCTX_ENABLE build must not fall back to WARN"; @@ -140,8 +116,7 @@ class RoctxNvmlRoctxMicrotest : public ::testing::Test { #endif } - // After a range's scope ends: exactly one enter and one exit, checked separately so a - // destructor that pushed instead of popping cannot pass as two events. + // One enter and one exit, counted apart so a destructor that pushes can't pass. static void ExpectOneRangeEnteredAndExited() { #ifdef ROCTX_ENABLE EXPECT_TRUE(g_logged.empty()) << "a ROCTX_ENABLE build must not fall back to WARN"; @@ -225,9 +200,7 @@ struct MixedPayload { } // namespace -// --------------------------------------------------------------------------- // Lookup tables -// --------------------------------------------------------------------------- TEST_F(RoctxNvmlRoctxMicrotest, NvtxToRoctx_MapsEachSupportedNvtxTypeToItsRoctxType) { EXPECT_EQ(nvtxToRoctx.size(), 4u); @@ -248,8 +221,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, NameTables_MatchTheEnumOrder) { for (int i = 0; i < ncclNumTypes; ++i) EXPECT_STREQ(ncclDataTypeStr[i], dataTypes[i]) << "type " << i; } -// One slot per roctx entry type. The DATATYPE slot has no initialiser today, but nothing in the -// tree reads this table, so that gap is left unpinned rather than locked in. +// One slot per roctx entry type (DATATYPE has no initialiser; nothing reads it). TEST_F(RoctxNvmlRoctxMicrotest, RoctxEntryTypeStr_HasOneSlotPerEntryType) { static_assert(sizeof(roctxEntryTypeStr) / sizeof(roctxEntryTypeStr[0]) == ROCTX_PAYLOAD_NUM_ENTRY_TYPES, "roctxEntryTypeStr must have one slot per roctxPayloadEntryType"); @@ -258,9 +230,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, RoctxEntryTypeStr_HasOneSlotPerEntryType) { EXPECT_STREQ(roctxEntryTypeStr[ROCTX_PAYLOAD_ENTRY_TYPE_REDOP], "ROCTX_PAYLOAD_ENTRY_TYPE_REDOP"); } -// --------------------------------------------------------------------------- // roctxAlloc / roctxFree -// --------------------------------------------------------------------------- TEST_F(RoctxNvmlRoctxMicrotest, Alloc_AllocatesEntriesAndMessage_FreeReleasesBoth) { roctxPayloadInfo info; @@ -284,9 +254,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, Alloc_EntriesAllocationFails_LeavesBothPointersN EXPECT_EQ(*calls, 1); } -// The message allocation fails after the entries succeeded: the entries buffer is released -// (under ASan, LeakSanitizer catches a missing free) and nulled, so a later roctxFree cannot -// double-free. +// The message allocation fails after the entries succeeded: entries are freed and nulled. TEST_F(RoctxNvmlRoctxMicrotest, Alloc_MessageAllocationFails_ReleasesEntriesAndNullsBoth) { int* calls = FailMallocAfter(1); roctxPayloadInfo info; @@ -304,9 +272,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, Free_ToleratesNullPointers) { EXPECT_EQ(info.message, nullptr); } -// --------------------------------------------------------------------------- // extractPayloadInfo / stringify -// --------------------------------------------------------------------------- TEST_F(RoctxNvmlRoctxMicrotest, Extract_NoEntriesBuffer_ReturnsWithoutTouchingInfo) { roctxPayloadInfo info; @@ -357,8 +323,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, Stringify_ZeroEntriesIsJustTheId) { EXPECT_EQ(Stringify("Empty", {}), "{Empty: }"); } -// Out-of-range redop/datatype values print "unknown" instead of indexing past the tables; -// the highest valid values sit right next to them so an off-by-one in either bound shows. +// Out-of-range values print "unknown"; the top valid values sit next to them to pin the bounds. TEST_F(RoctxNvmlRoctxMicrotest, Stringify_OutOfRangeRedOpAndDataTypePrintUnknown) { std::vector entries(4); entries[0] = {"lastOp", ROCTX_PAYLOAD_ENTRY_TYPE_REDOP, {}}; @@ -385,8 +350,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, Stringify_MissingBufferIsANoOp) { EXPECT_EQ(info.message, nullptr); } -// Entries present but no message buffer: still a no-op, checked separately from the -// both-null case above because the guard is `!entries || !message`. +// Entries but no message: still a no-op (the guard is `!entries || !message`). TEST_F(RoctxNvmlRoctxMicrotest, Stringify_EntriesWithoutMessageIsANoOp) { roctxPayloadSchemaEntryInfo entry{"n", ROCTX_PAYLOAD_ENTRY_TYPE_INT, {}}; roctxPayloadInfo info; @@ -397,8 +361,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, Stringify_EntriesWithoutMessageIsANoOp) { EXPECT_EQ(info.message, nullptr); } -// A mapping to a roctx type outside the enum reaches extractPayloadInfo's default branch, -// which stores the type but no payload, and stringify then prints the unknown-type text. +// An out-of-enum roctx type stores no payload and prints the unknown-type text. TEST_F(RoctxNvmlRoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefaultBranch) { const uint64_t nvtxType = NVTX_PAYLOAD_ENTRY_TYPE_INT + 2; ASSERT_EQ(nvtxToRoctx.count(nvtxType), 0u); @@ -415,8 +378,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefault roctxFree(&info); } -// Pinned as observed: an nvtx type with no mapping is looked up with std::map::operator[], -// which inserts it as ROCTX_PAYLOAD_ENTRY_TYPE_INT (value 0) and decodes the field as an int. +// Pinned as observed: an unmapped nvtx type is inserted by operator[] as INT and decoded as one. TEST_F(RoctxNvmlRoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInsertedIntoTheMap) { const uint64_t unmapped = NVTX_PAYLOAD_ENTRY_TYPE_INT + 1; ASSERT_EQ(nvtxToRoctx.count(unmapped), 0u); @@ -432,9 +394,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInserte roctxFree(&info); } -// --------------------------------------------------------------------------- // roctx_scoped_range_in -// --------------------------------------------------------------------------- TEST_F(RoctxNvmlRoctxMicrotest, LogRoctxParam_IsReadFromRcclLogRoctxWithDefaultZero) { std::string seenEnv; @@ -459,8 +419,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_LoggingOff_DoesNothing) { EXPECT_TRUE(g_logged.empty()); } -// One range is one event on enter and one on exit: a roctracer push/pop with ROCTX_ENABLE, -// otherwise the WARN that stands in for each. +// One range: one enter event and one exit event. TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_LoggingOn_OneEventOnEnterAndOneOnExit) { SetLogRoctx(1); { @@ -480,10 +439,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_DefaultConstructorDelegatesToTheMess ExpectOneRangeEnteredAndExited(); } -// The payload constructor builds the message from the schema (pushed verbatim with ROCTX_ENABLE) -// and the destructor frees it in both builds; under ASan, LeakSanitizer flags a missing roctxFree. -// The two mallocs (entries, message) are checked directly too, since without ROCTX_ENABLE the -// pushed message is not observable. +// The constructor pushes the schema-built message; the destructor frees it in both builds. TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOn_PushesTheStringifiedPayload) { SetLogRoctx(1); int* mallocs = CountMallocs(); @@ -499,8 +455,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOn_Push ExpectOneRangeEnteredAndExited(); } -// roctxAlloc fails inside the constructor, leaving the message null: construction and -// destruction must still complete without crashing and report one enter and one exit. +// roctxAlloc fails in the constructor: the range still enters and exits once, without crashing. TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOnAndAllocFailure_StillEntersAndExits) { SetLogRoctx(1); int* mallocs = FailMallocAfter(0); @@ -524,9 +479,7 @@ TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_PayloadConstructorWithLoggingOff_Doe EXPECT_TRUE(g_logged.empty()); } -// --------------------------------------------------------------------------- // nvmlwrap_stub.cc: no NVML on ROCm, so every query reports "no device". -// --------------------------------------------------------------------------- namespace { nvmlDevice_t FakeDevice() { return reinterpret_cast(0x1234); } @@ -589,8 +542,7 @@ TEST(RoctxNvmlStubMicrotest, DeviceLookupsFailWithoutWritingOutputs) { EXPECT_EQ(value, 99u); } -// These fail too, but clear their output first so a caller that ignores the error -// reads a defined value instead of stack garbage. +// These fail too, but clear their output first. TEST(RoctxNvmlStubMicrotest, FailingQueriesClearTheirOutputs) { char name[8] = "garbage"; EXPECT_EQ(ncclNvmlDeviceGetName(FakeDevice(), name, sizeof(name)), ncclSystemError); From f09b20b1fe2eec1a600af3ffca94bf54cffbb762 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 15:40:46 -0700 Subject: [PATCH 12/16] test(rccl): trim microtest docs and comments to match the other host tests Shorten the MICROTEST_README entry, CMake comments and test-file comments to the length the other host microtests use. Comment and doc changes only. Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 7 ++-- projects/rccl/test/host/MICROTEST_README.md | 18 ++-------- projects/rccl/test/host/rocmwrap-test.cc | 40 ++++----------------- 3 files changed, 11 insertions(+), 54 deletions(-) diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 88e618ffd12f..40478f36372e 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -389,11 +389,8 @@ function(rccl_define_micro_source_lists) ../common/main_altrsmi.cpp PARENT_SCOPE) - # rccl-UnitTestsMicroRocmwrap -- src/misc/rocmwrap.cc. Its own binary: fakes/rocmwrap_fakes.cc - # and fakes/nccl_fakes.cc (linked by most other targets) define rocmLibraryInit, - # ncclCuMemEnable and ncclCuMemHandleType, which this target compiles for real. The test TU - # supplies ncclDebugLog, g_loadParam and the HSA entry points itself; kernel_config.cc and - # archinfo.cc are the real thing, not fakes. + # rccl-UnitTestsMicroRocmwrap -- src/misc/rocmwrap.cc. Own binary: fakes/rocmwrap_fakes.cc + # and fakes/nccl_fakes.cc define rocmLibraryInit, ncclCuMemEnable and ncclCuMemHandleType. set(TEST_MICRO_ROCMWRAP_SOURCE_FILES rocmwrap-test.cc fakes/hip_fakes.cc diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index e8bd2dba986c..c096e3327395 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -311,21 +311,9 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: sharing `rccl-UnitTestsMicro`: `gin-plugin-init-test.cc` already defines `ncclParamGinEnable` there. See `test_categories_micro_gin_host.yaml`. - **`rccl-UnitTestsMicroRocmwrap`** — `src/misc/rocmwrap.cc` (via - `ROCMWRAP_CC_PATH`, from `rocmwrap-test.cc`); suite `RocmwrapMicrotest.*`. - Covers the kernel-version probe, the cuMem VMM functional probe and its - cleanup, the capability gates behind `ncclIsCuMemSupported` / - `ncclCuMemRuntimeSupported`, and `rocmLibraryInit`'s ROCr version and - DMA-BUF setup. Its own binary because `fakes/rocmwrap_fakes.cc` and - `fakes/nccl_fakes.cc` define `rocmLibraryInit`, `ncclCuMemEnable` and - `ncclCuMemHandleType`. The test TU defines `ncclDebugLog`, `g_loadParam`, - `hsa_system_get_info` and `hsa_init`, and seams `uname`, `fopen`, - `hipDriverGetVersion`, `hipDeviceGet`, the kernel-config reader and the weak - `hsa_amd_portable_export_dmabuf` symbol inside `rocmwrap.cc` only; - `kernel_config.cc` and `archinfo.cc` are compiled for real. The in-window - arm of `ncclCuMemEnable` compiles only on ROCm 7.12+ or the 7.0.2.x backport - (`NCCL_CUMEM_VERSION_SUPPORTED`), and that of `ncclCuMemHostEnable` only on - 7.12+ (`NCCL_CUMEM_HOST_VERSION_SUPPORTED`), so on 7.1/7.2 the tests pin the - out-of-window arms. See + `ROCMWRAP_CC_PATH`, from `rocmwrap-test.cc`); suite `RocmwrapMicrotest.*`. Its + own binary because `fakes/rocmwrap_fakes.cc` and `fakes/nccl_fakes.cc` define + `rocmLibraryInit`, `ncclCuMemEnable` and `ncclCuMemHandleType`. See `test_categories_micro_rocmwrap.yaml`. - **`rccl-UnitTestsMicroDiagnostics`**: `src/diagnostics/p2p.cc` (via `DIAG_P2P_CC_PATH`, suite `DiagP2pMicrotest.*`) and `src/diagnostics/ib_write_bw.cc` diff --git a/projects/rccl/test/host/rocmwrap-test.cc b/projects/rccl/test/host/rocmwrap-test.cc index aaf066d69b78..8bb7e5c8ba06 100644 --- a/projects/rccl/test/host/rocmwrap-test.cc +++ b/projects/rccl/test/host/rocmwrap-test.cc @@ -4,25 +4,9 @@ * See LICENSE.txt for license information ************************************************************************/ -// Host-only microtests for src/misc/rocmwrap.cc: the kernel-version probe, the -// cuMem VMM functional probe and capability gates, the NCCL_CUMEM_ENABLE / -// cuMem-host decisions, and rocmLibraryInit's one-time ROCr + DMA-BUF setup. -// -// The real HIP entry points are the shared fakes/hip_fakes.cc hooks. Everything -// else rocmwrap.cc reaches outside HIP goes through a seam scoped to its own -// function bodies (see the #define block before ROCMWRAP_CC_PATH): uname, fopen -// (/proc/kallsyms), hipDriverGetVersion and hipDeviceGet (fixed-value stubs in -// hip_fakes.cc), the kernel-config reader, and the weakly linked -// hsa_amd_portable_export_dmabuf entry. hsa_system_get_info and hsa_init are -// defined here. kernel_config.cc and archinfo.cc are compiled for real, so the -// option matching and the gfx-arch parsing are the production code. -// -// ncclCuMemEnable and ncclCuMemHostEnable each have an arm gated at compile -// time on HIP_VERSION: ncclCuMemEnable's on the cuMem window (7.12.60540, or the -// 7.0.2.x backport), ncclCuMemHostEnable's on 7.12.60540 only. ROCm 7.1/7.2 -// (local and CI) are outside both, so those tests pin the out-of-window arms; -// the host arm also needs hipDeviceAttributeHostNumaId, which these headers do -// not declare. +// Host-only microtests for src/misc/rocmwrap.cc. Non-HIP calls are seamed inside rocmwrap.cc only +// (see the #defines before ROCMWRAP_CC_PATH). On ROCm 7.1/7.2 the cuMem tests pin the +// out-of-window arms of ncclCuMemEnable/ncclCuMemHostEnable. #include @@ -51,8 +35,7 @@ #include "hsa/hsa.h" #include "bootstrap.h" -// --- Logging and params: this binary links no shared nccl fakes. ------------- - +// --- Logging and params: this binary links no shared nccl fakes. uint32_t ncclDebugLevelMask = NCCL_DEBUG_LEVEL_MASK_UNINITIALIZED; // log everything uint64_t ncclDebugMask = ~uint64_t(0); @@ -82,8 +65,7 @@ std::function g_loadParam = [](const char* env, i #include "fakes/param_redirect.h" -// --- HSA: the three ROCr entry points rocmwrap.cc calls. ---------------------- - +// --- HSA: the three ROCr entry points rocmwrap.cc calls. namespace { std::function g_hsaSystemGetInfo; int g_hsaInitCalls = 0; @@ -97,7 +79,7 @@ extern "C" hsa_status_t hsa_init() { return HSA_STATUS_SUCCESS; } -// --- Seams scoped to rocmwrap.cc's own function bodies. ----------------------- +// --- Seams scoped to rocmwrap.cc's own function bodies. // Every header rocmwrap.cc includes is pulled in above, so their include guards // keep the renames below from rewriting any declaration. @@ -313,9 +295,7 @@ class RocmwrapMicrotest : public ::testing::Test { } // namespace -// --------------------------------------------------------------------------- // ncclGetKernelVersionCode -// --------------------------------------------------------------------------- TEST_F(RocmwrapMicrotest, KernelVersionCode_PacksMajorAndMinorFromTheRelease) { SetKernelRelease("6.11.4-45-generic"); @@ -334,9 +314,7 @@ TEST_F(RocmwrapMicrotest, KernelVersionCode_UnparseableReleaseIsZero) { EXPECT_EQ(ncclGetKernelVersionCode(), 0); } -// --------------------------------------------------------------------------- // ncclCuMemFunctionalProbe -// --------------------------------------------------------------------------- TEST_F(RocmwrapMicrotest, FunctionalProbe_AllStepsSucceed_ReturnsOneAndUndoesEveryStep) { VmmLog log; @@ -409,9 +387,7 @@ TEST_F(RocmwrapMicrotest, FunctionalProbe_FailingCleanupIsReportedButDoesNotChan EXPECT_TRUE(Logged(NCCL_LOG_INFO, "HIP failure")); } -// --------------------------------------------------------------------------- // ncclCuMemCapabilityCheck / ncclIsCuMemSupported -// --------------------------------------------------------------------------- TEST_F(RocmwrapMicrotest, CapabilityCheck_AllGatesPassAndProbeSucceeds_ReturnsOne) { VmmLog log; @@ -566,9 +542,7 @@ TEST_F(RocmwrapMicrotest, RuntimeSupported_SkipsTheArchGateAndRunsTheCheckOnce) ::testing::ExitedWithCode(0), ""); } -// --------------------------------------------------------------------------- // ncclCuMemEnable / ncclCuMemHostEnable (out-of-window HIP_VERSION arm) -// --------------------------------------------------------------------------- #if !NCCL_CUMEM_VERSION_SUPPORTED(HIP_VERSION) TEST_F(RocmwrapMicrotest, CuMemEnable_OutsideTheCompileTimeWindowIsAlwaysZero) { @@ -597,9 +571,7 @@ TEST_F(RocmwrapMicrotest, CuMemHostEnable_ReturnsTheCachedValue) { EXPECT_EQ(ncclCuMemHostEnable(), 1); } -// --------------------------------------------------------------------------- // initOnceFunc / rocmLibraryInit -// --------------------------------------------------------------------------- // CUDA_LAUNCH_BLOCKING is on for any non-empty value other than exactly "0". TEST_F(RocmwrapMicrotest, Init_ParsesCudaLaunchBlocking) { From 6da3065a74cf7a3e8abf9d49340484c6acc2ea40 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 15:40:59 -0700 Subject: [PATCH 13/16] test(rccl): trim microtest docs and comments to match the other host tests Shorten the MICROTEST_README entry, CMake comments and test-file comments to the length the other host microtests use. Comment and doc changes only. Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 6 +-- projects/rccl/test/host/MICROTEST_README.md | 12 +---- projects/rccl/test/host/amdsmi-wrap-test.cc | 54 +++------------------ 3 files changed, 10 insertions(+), 62 deletions(-) diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 6ea508320e4b..85b365c79876 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -373,10 +373,8 @@ function(rccl_define_micro_source_lists) ../common/main_altrsmi.cpp PARENT_SCOPE) - # rccl-UnitTestsMicroAmdsmiWrap -- src/misc/amdsmi_wrap.cc. Its own binary: fakes/amdsmi_fakes.cc, - # linked by most other targets, defines the amd_smi_* wrappers this target compiles for real. The - # test TU supplies ncclDebugLog, g_loadParam and the ARSMI_* (alt_rsmi.cc) stand-ins itself; - # utils.cc is the real thing for busIdToInt64. + # rccl-UnitTestsMicroAmdsmiWrap -- src/misc/amdsmi_wrap.cc. Own binary: fakes/amdsmi_fakes.cc + # defines the amd_smi_* wrappers this target compiles for real. set(TEST_MICRO_AMDSMI_WRAP_SOURCE_FILES amdsmi-wrap-test.cc fakes/hip_fakes.cc diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index d5ec35230aa8..fe8cf10cbdb0 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -312,16 +312,8 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: `ncclParamGinEnable` there. See `test_categories_micro_gin_host.yaml`. - **`rccl-UnitTestsMicroAmdsmiWrap`** — `src/misc/amdsmi_wrap.cc` (via `AMDSMI_WRAP_CC_PATH`, from `amdsmi-wrap-test.cc`); suite - `AmdsmiWrapMicrotest.*`. Covers init/shutdown and their caching, GPU count, - PCI bus-id lookups both ways, link info, firmware version, fabric discovery - (including the amd_smi fabric ABI probe and its sysfs fallbacks) and - telemetry, each across the WSL2, amd_smi_lib and ARSMI backends. Its own - binary because `fakes/amdsmi_fakes.cc`, linked by most other targets, - defines the `amd_smi_*` wrappers this file compiles for real. The amd_smi - library is a scripted `FakeSmi` installed through the file's own `pfn_*` - table; `ARSMI_*` (`alt_rsmi.cc`) is defined in the test TU, and `access`, - `dlopen`/`dlsym` and `hipDeviceGetByPCIBusId` are seamed inside - `amdsmi_wrap.cc` only; `utils.cc` is real for `busIdToInt64`. See + `AmdsmiWrapMicrotest.*`. Its own binary because `fakes/amdsmi_fakes.cc` + defines the `amd_smi_*` wrappers this file compiles. See `test_categories_micro_amdsmi_wrap.yaml`. - **`rccl-UnitTestsMicroDiagnostics`**: `src/diagnostics/p2p.cc` (via `DIAG_P2P_CC_PATH`, suite `DiagP2pMicrotest.*`) and `src/diagnostics/ib_write_bw.cc` diff --git a/projects/rccl/test/host/amdsmi-wrap-test.cc b/projects/rccl/test/host/amdsmi-wrap-test.cc index 44392722462d..961af1978b96 100644 --- a/projects/rccl/test/host/amdsmi-wrap-test.cc +++ b/projects/rccl/test/host/amdsmi-wrap-test.cc @@ -4,28 +4,9 @@ * See LICENSE.txt for license information ************************************************************************/ -// Host-only microtests for src/misc/amdsmi_wrap.cc: RCCL's wrapper over GPU -// enumeration, PCI bus IDs, link topology, firmware version and UALoE/UALink -// fabric discovery. Every query has three backends: WSL2 (HIP calls only), -// amd_smi_lib (dlopen'd, opted into with RCCL_USE_AMD_SMI_LIB=1) and the -// internal ARSMI sysfs reader (the default). -// -// The amd_smi_lib backend is reached through the file's own pfn_* table, which -// this TU can assign because it #includes the .cc; FakeSmi below stands in for -// the library as a list of GPUs. ARSMI_* (src/misc/alt_rsmi.cc) is defined -// here as scriptable stand-ins. access (the WSL2 probe), dlopen/dlsym and -// hipDeviceGetByPCIBusId are seamed inside amdsmi_wrap.cc's own bodies only; -// the other HIP calls go through fakes/hip_fakes.cc, and utils.cc's -// busIdToInt64 is the real code. -// -// The file has process-wide state the tests reset by hand (see Reset()): the -// init/fabric caches, the WSL2 flag, and the dlopen handle. The handle is a -// function-local static, so the two dlopen tests run in forked children. -// -// The fabric ABI probe compares what the loaded library writes against the -// layout this build was compiled for. FakeSmi writes the compiled layout's -// extent by default, so the typed path is taken whatever amd_smi header the -// build found; the mismatch tests write the other extents. +// Host-only microtests for src/misc/amdsmi_wrap.cc across its WSL2, amd_smi_lib and ARSMI +// backends. FakeSmi stands in for amd_smi_lib through the file's own pfn_* table; Reset() clears +// the file's process-wide caches, and the dlopen tests fork because the handle is static. #include @@ -53,8 +34,7 @@ #include "core.h" #include "utils.h" -// --- Logging and params: this binary links no shared nccl fakes. ------------- - +// --- Logging and params: this binary links no shared nccl fakes. uint32_t ncclDebugLevelMask = NCCL_DEBUG_LEVEL_MASK_UNINITIALIZED; // log everything uint64_t ncclDebugMask = ~uint64_t(0); thread_local int ncclDebugNoWarn = 0; // utils.cc reads it @@ -85,8 +65,7 @@ std::function g_loadParam = [](const char* env, i #include "fakes/param_redirect.h" -// --- ARSMI (src/misc/alt_rsmi.cc): scriptable stand-ins. ---------------------- - +// --- ARSMI (src/misc/alt_rsmi.cc): scriptable stand-ins. namespace { struct FakeArsmi { int initResult = 0; @@ -133,8 +112,7 @@ int ARSMI_get_fabric_info(uint32_t i, struct ARSMI_fabricInfo* info) { } const char* ARSMI_fabric_telem_id_to_string(uint64_t) { return "ARSMI_NAME"; } -// --- Seams scoped to amdsmi_wrap.cc's own function bodies. -------------------- - +// --- Seams scoped to amdsmi_wrap.cc's own function bodies. namespace { std::function g_access; std::function g_dlopen; @@ -499,9 +477,7 @@ class AmdsmiWrapMicrotest : public ::testing::Test { } // namespace -// --------------------------------------------------------------------------- // amd_smi_init / amd_smi_shutdown -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, Init_Wsl2SkipsEveryBackend) { UseWsl2(); @@ -644,9 +620,7 @@ TEST_F(AmdsmiWrapMicrotest, Init_DlopenResolvesEveryEntryPointByName) { EXPECT_EXIT(std::exit(DlopenResolvesChild()), ::testing::ExitedWithCode(0), ""); } -// --------------------------------------------------------------------------- // getProcessorHandle (via amd_smi_getFirmwareVersion, its only plain caller) -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_IsAmdSmiLibOnly) { amdsmi_processor_handle h = nullptr; @@ -701,9 +675,7 @@ TEST_F(AmdsmiWrapMicrotest, ProcessorHandle_MissingSymbolFailsQuietly) { } #endif -// --------------------------------------------------------------------------- // amd_smi_getNumDevice -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, NumDevice_Wsl2AsksHip) { UseWsl2(); @@ -745,9 +717,7 @@ TEST_F(AmdsmiWrapMicrotest, NumDevice_ArsmiReportsItsCount) { EXPECT_EQ(amd_smi_getNumDevice(&n), ncclInternalError); } -// --------------------------------------------------------------------------- // amd_smi_getDevicePciBusIdString -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, PciBusIdString_Wsl2AsksHip) { UseWsl2(); @@ -807,9 +777,7 @@ TEST_F(AmdsmiWrapMicrotest, PciBusIdString_HonorsTheBufferLength) { EXPECT_STREQ(busId, "002a:"); } -// --------------------------------------------------------------------------- // amd_smi_getDeviceIndexByPciBusId -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_Wsl2AsksHip) { UseWsl2(); @@ -900,9 +868,7 @@ TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_LibLookupFailureIsInternalError) { EXPECT_EQ(amd_smi_getDeviceIndexByPciBusId("0000:c3:00.0", &idx), ncclInternalError); } -// --------------------------------------------------------------------------- // amd_smi_getLinkInfo -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, LinkInfo_Wsl2IsOnePcieHop) { UseWsl2(); @@ -1029,9 +995,7 @@ TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibUnknownIndexIsLogged) { EXPECT_EQ(g_smi.lastLinkDst, nullptr); } -// --------------------------------------------------------------------------- // amd_smi_getFirmwareVersion -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_InitFailureIsReturned) { g_arsmi.initResult = 1; @@ -1089,9 +1053,7 @@ TEST_F(AmdsmiWrapMicrotest, FirmwareVersion_LibFailures) { EXPECT_TRUE(Logged(NCCL_LOG_ERROR, "amdsmi_get_fw_info symbol not loaded")); } -// --------------------------------------------------------------------------- // amd_smi_ensureFabricInitialized and the fabric accessors -// --------------------------------------------------------------------------- namespace { ARSMI_fabricInfo ArsmiFabric(int supported, uint32_t accelId, uint32_t bw) { @@ -1329,9 +1291,7 @@ TEST_F(AmdsmiWrapMicrotest, FabricAccessors_PropagateAnInitFailure) { EXPECT_EQ(amd_smi_getFabricBandwidth(0, &bw), ncclInternalError); } -// --------------------------------------------------------------------------- // Fabric telemetry -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, Telemetry_IsAmdSmiLibOnly) { amdsmi_fabric_telemetry_t* t = nullptr; @@ -1380,9 +1340,7 @@ TEST_F(AmdsmiWrapMicrotest, Telemetry_FailuresAreReturned) { EXPECT_EQ(amd_smi_freeFabricTelemetry(0, dummy), ncclInternalError); } -// --------------------------------------------------------------------------- // amd_smi_fabricTelemIdToString -// --------------------------------------------------------------------------- TEST_F(AmdsmiWrapMicrotest, TelemIdToString_WithoutTheSymbolAsksArsmi) { pfn_amdsmi_fabric_telem_id_to_string = nullptr; From 521c4ce7de18e47f9f844f180b3a4a81f558f4ce Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Wed, 14 Oct 2026 16:59:05 -0700 Subject: [PATCH 14/16] test(rccl): follow the microtest naming and registration rules Give the binary its own test_categories_micro_.yaml and register it in ci-precheckin.json and test_executor.py's llvm-cov list, as MICROTEST_README's naming section requires, so the CI test runner and coverage pick it up. Co-Authored-By: Claude --- projects/rccl/test/host/CMakeLists.txt | 2 +- projects/rccl/test/host/MICROTEST_README.md | 3 +- projects/rccl/test/host/run_host_tests.sh | 3 +- projects/rccl/test/test_categories_micro.yaml | 2 - .../test_categories_micro_roctx_nvml.yaml | 46 +++++++++++++++++++ .../test_runner/configs/ci-precheckin.json | 17 +++++++ .../scripts/test_runner/lib/test_executor.py | 3 +- 7 files changed, 69 insertions(+), 7 deletions(-) create mode 100644 projects/rccl/test/test_categories_micro_roctx_nvml.yaml diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 82f11ad412aa..ac27a3f54dc4 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -1074,7 +1074,7 @@ set(_roctx_nvml_common_args ${_MICRO_MODE_DEFS} ROCTX_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/roctx.cc" NVMLWRAP_STUB_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/nvmlwrap_stub.cc" - CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml") + CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_roctx_nvml.yaml") rccl_add_micro_binary(rccl-UnitTestsMicroRoctxNvml ${_roctx_nvml_common_args}) # -roctx builds roctx.cc's ROCTX_ENABLE arm against the stand-in fakes/roctx_include. diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index 295e8b70ffa3..1dabe189afbd 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -251,7 +251,8 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: - **`rccl-UnitTestsMicroRoctxNvml`** (+ **`-roctx`**) — `src/misc/roctx.cc` and `src/misc/nvmlwrap_stub.cc` (via `ROCTX_CC_PATH` / `NVMLWRAP_STUB_CC_PATH`, from `roctx-nvml-stub-test.cc`); suites `RoctxNvml*`. Its own binary because - `fakes/comm_fakes.cc` stubs the `roctx_scoped_range_in` this file defines. + `fakes/comm_fakes.cc` stubs the `roctx_scoped_range_in` this file defines. See + `test_categories_micro_roctx_nvml.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.*`, diff --git a/projects/rccl/test/host/run_host_tests.sh b/projects/rccl/test/host/run_host_tests.sh index 4591a0bf5f14..7fd9496dbd85 100755 --- a/projects/rccl/test/host/run_host_tests.sh +++ b/projects/rccl/test/host/run_host_tests.sh @@ -168,8 +168,7 @@ do_host_tests() { "rccl-UnitTestsMicroCe:$SCRIPT_DIR/host_tests_micro_ce.xml" "rccl-UnitTestsMicroClientSupport:$SCRIPT_DIR/host_tests_micro_client_support.xml" "rccl-UnitTestsMicroRoctxNvml:$SCRIPT_DIR/host_tests_micro_roctx_nvml.xml" - # roctx.cc gates its roctracer calls on ROCTX_ENABLE at the preprocessor; this - # variant builds that arm against a stand-in roctracer header. See test/host/CMakeLists.txt. + # Standalone-only: roctx.cc's ROCTX_ENABLE arm. "rccl-UnitTestsMicroRoctxNvml-roctx:$SCRIPT_DIR/host_tests_micro_roctx_nvml_roctx.xml" "rccl-UnitTestsMicroGinHost:$SCRIPT_DIR/host_tests_micro_gin_host.xml" "rccl-UnitTestsMicroDiagnostics:$SCRIPT_DIR/host_tests_micro_diagnostics.xml" diff --git a/projects/rccl/test/test_categories_micro.yaml b/projects/rccl/test/test_categories_micro.yaml index 9675aa4eb966..a940ce5f4f7c 100644 --- a/projects/rccl/test/test_categories_micro.yaml +++ b/projects/rccl/test/test_categories_micro.yaml @@ -114,8 +114,6 @@ test_categories: # ce_coll.cc (rccl-UnitTestsMicroCe) - "CeCollMicrotest.*" - "LaunchPaths/CeAlltoAllvSyncMicrotest.*" - # misc/roctx.cc, misc/nvmlwrap_stub.cc (rccl-UnitTestsMicroRoctxNvml[-roctx]) - - "RoctxNvml*" # dev_runtime.cc - "Alloc*" - "Comm*" diff --git a/projects/rccl/test/test_categories_micro_roctx_nvml.yaml b/projects/rccl/test/test_categories_micro_roctx_nvml.yaml new file mode 100644 index 000000000000..b628c37ead26 --- /dev/null +++ b/projects/rccl/test/test_categories_micro_roctx_nvml.yaml @@ -0,0 +1,46 @@ +# rccl-UnitTestsMicroRoctxNvml — shared/ctest schema. +# Host-only micro-tests for misc/roctx.cc and misc/nvmlwrap_stub.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: µ_roctx_nvml_quick + description: "Smoke — host-only misc/roctx.cc and misc/nvmlwrap_stub.cc micro-tests, sub-second." + test_patterns: + - "RoctxNvml*" + exclude_windows: + - "*" + labels: + - "rccl" + - "quick" + - "smoke" + - "precheck" + - "microtest" + standard: + <<: *micro_roctx_nvml_quick + labels: + - "rccl" + - "standard" + - "regression" + - "microtest" + comprehensive: + <<: *micro_roctx_nvml_quick + labels: + - "rccl" + - "comprehensive" + - "comprehensive_regression" + - "microtest" + full: + <<: *micro_roctx_nvml_quick + labels: + - "rccl" + - "full" + - "full_regression" + - "microtest" + +execution_settings: + timeout_multiplier: 1 + category_timeouts: + quick: 600 + standard: 600 + comprehensive: 600 + full: 600 diff --git a/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json b/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json index 07da6b28ba73..bb686e5cf5c4 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json +++ b/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json @@ -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_roctx_nvml": { + "is_gtest": true, + "binary": "host/rccl-UnitTestsMicroRoctxNvml", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 60, + "tests": [ + {"name": "RoctxNvmlMicrotests", "description": "Host-only misc/roctx.cc + misc/nvmlwrap_stub.cc microtests (no GPU, no HIP/RCCL runtime)", "test_filter": "*"} + ] } }, "test_suites": [ @@ -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 - misc/roctx.cc + misc/nvmlwrap_stub.cc", + "description": "rccl-UnitTestsMicroRoctxNvml (host-only misc/roctx.cc + misc/nvmlwrap_stub.cc microtests, no GPU/HIP/RCCL runtime)", + "config": "unit_tests_micro_roctx_nvml", + "enabled": true } ] } diff --git a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py index 70add7053ab2..cc6d573608ee 100644 --- a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py +++ b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py @@ -3292,7 +3292,8 @@ def generate_coverage_report(self): "rccl-UnitTestsMicroInit-faultinj", "rccl-UnitTestsMicroEnqueue", "rccl-UnitTestsMicroSymKernels", - "rccl-UnitTestsMicroGinHost"]: + "rccl-UnitTestsMicroGinHost", + "rccl-UnitTestsMicroRoctxNvml"]: binary_path = os.path.join(host_test_dir, binary) if os.path.isfile(binary_path): object_files.extend(["--object", binary_path]) From 74e100d3edd059b52f8163981dd8840853114e44 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Thu, 15 Oct 2026 09:36:29 -0700 Subject: [PATCH 15/16] test(rccl): pin amdsmi_wrap's link-info type check and fabric bounds Add a hip-id-1 CPU after GPU1 so the link lookup needs its GPU type check, and populate the out-of-range fabric slot so the bounds checks are needed. Co-Authored-By: Claude --- projects/rccl/test/host/amdsmi-wrap-test.cc | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/projects/rccl/test/host/amdsmi-wrap-test.cc b/projects/rccl/test/host/amdsmi-wrap-test.cc index 961af1978b96..0fedc4659d5b 100644 --- a/projects/rccl/test/host/amdsmi-wrap-test.cc +++ b/projects/rccl/test/host/amdsmi-wrap-test.cc @@ -953,6 +953,10 @@ TEST_F(AmdsmiWrapMicrotest, LinkInfo_ArsmiFailureIsInternalError) { TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibQueriesTheSrcAndDstGpus) { InitFor(/*lib=*/true); TwoGpusBehindACpu(); + FakeProc trailingCpu; // after GPU1 too, so a last-match lookup without the type check picks it + trailingCpu.gpu = false; + trailingCpu.hipId = 1; + g_smi.sockets[0].push_back(trailingCpu); g_smi.linkType = AMDSMI_LINK_TYPE_PCIE; g_smi.linkHops = 3; g_smi.minBw = 64; // only XGMI uses the bandwidth ratio, so a PCIe link's count stays 1 @@ -1262,6 +1266,9 @@ TEST_F(AmdsmiWrapMicrotest, FabricAccessors_ReadTheCacheWithBoundsChecks) { EXPECT_TRUE(supported); ASSERT_EQ(amd_smi_isFabricSupported(1, &supported), ncclSuccess); EXPECT_FALSE(supported); + // Past the device count but populated, so only the bounds checks keep index 2 out. + amdsmiFabricDevices[2].fabricSupported = true; + amdsmiFabricDevices[2].bandwidth = 500000; supported = true; ASSERT_EQ(amd_smi_isFabricSupported(2, &supported), ncclSuccess); EXPECT_FALSE(supported); From 48a7554ea53490c655364eed3ed003f72813bd13 Mon Sep 17 00:00:00 2001 From: Alex Duforest Date: Fri, 16 Oct 2026 14:26:21 -0700 Subject: [PATCH 16/16] test(rccl): name the roctx/rocmwrap/amdsmi pins KnownBug_ Also reset g_legacyTelemName in Reset() and drop the Jira key from the roctx-nvml-stub-test.cc header. Co-Authored-By: Claude --- projects/rccl/test/host/amdsmi-wrap-test.cc | 18 +++++++++--------- projects/rccl/test/host/rocmwrap-test.cc | 8 ++++---- .../rccl/test/host/roctx-nvml-stub-test.cc | 6 +++--- 3 files changed, 16 insertions(+), 16 deletions(-) diff --git a/projects/rccl/test/host/amdsmi-wrap-test.cc b/projects/rccl/test/host/amdsmi-wrap-test.cc index 0fedc4659d5b..8afd388a89f6 100644 --- a/projects/rccl/test/host/amdsmi-wrap-test.cc +++ b/projects/rccl/test/host/amdsmi-wrap-test.cc @@ -428,6 +428,7 @@ class AmdsmiWrapMicrotest : public ::testing::Test { g_dlopen = [](const char*, int) -> void* { return nullptr; }; g_dlsymTable.clear(); g_hipDeviceGetByPCIBusId = [](int*, const char*) { return hipErrorInvalidValue; }; + g_legacyTelemName = "LEGACY_NAME"; InstallFakeSmiTable(); is_wsl2 = -1; @@ -689,8 +690,8 @@ TEST_F(AmdsmiWrapMicrotest, NumDevice_Wsl2AsksHip) { EXPECT_EQ(n, 3u); } -// Counts every processor handle in every socket, CPUs included (pinned as observed). -TEST_F(AmdsmiWrapMicrotest, NumDevice_LibCountsProcessorsAcrossSockets) { +// Counts every processor handle in every socket, CPUs included; a fix flips this pin. +TEST_F(AmdsmiWrapMicrotest, NumDevice_KnownBug_LibCountsCpuProcessors) { InitFor(/*lib=*/true); g_smi.sockets = {{FakeProc{}, FakeProc{}}, {FakeProc{}}}; g_smi.sockets[1][0].gpu = false; @@ -731,8 +732,8 @@ TEST_F(AmdsmiWrapMicrotest, PciBusIdString_Wsl2AsksHip) { EXPECT_STREQ(busId, "wsl-2"); } -// Domain from bits [63:32], bus [15:8], device [7:3], function [2:0] (pinned as observed: the -// comment in the source describes the domain as [63:16]). +// Domain from bits [63:32], bus [15:8], device [7:3], function [2:0]; the source comment says +// the domain is [63:16]. TEST_F(AmdsmiWrapMicrotest, PciBusIdString_ArsmiFormatsTheBdf) { InitFor(/*lib=*/false); g_arsmi.bdfIds = {0, 0x0000002A0000C3DBull}; @@ -837,9 +838,8 @@ TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_LibPassesTheBdfFieldsAndReturnsTheHi EXPECT_EQ(g_smi.lastBdf.domain_number, 0x2au); } -// Pinned as observed, not endorsed: the device number is taken as `(busid & 0xff) >> 4`, which keeps -// only its low hex digit, so device 0x1b reaches amd_smi as 0x0b. Devices below 0x10 are unaffected. -TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_LibDropsTheDeviceNumbersHighDigit) { +// `(busid & 0xff) >> 4` keeps only the device number's low hex digit (0x1b becomes 0x0b); a fix flips this pin. +TEST_F(AmdsmiWrapMicrotest, IndexByPciBusId_KnownBug_LibDropsDeviceHighDigit) { InitFor(/*lib=*/true); g_smi.fromBdfStatus = AMDSMI_STATUS_NOT_FOUND; uint32_t idx = 0; @@ -987,8 +987,8 @@ TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibXgmiCountIsTheBandwidthRatio) { EXPECT_EQ(count, 1); } -// An unknown index is logged, and the query still runs against the null handle (pinned as observed). -TEST_F(AmdsmiWrapMicrotest, LinkInfo_LibUnknownIndexIsLogged) { +// An unknown index is logged, then queried as a null handle and reported as success; a fix flips this pin. +TEST_F(AmdsmiWrapMicrotest, LinkInfo_KnownBug_LibUnknownIndexStillSucceeds) { InitFor(/*lib=*/true); TwoGpusBehindACpu(); amdsmi_link_type_t type; diff --git a/projects/rccl/test/host/rocmwrap-test.cc b/projects/rccl/test/host/rocmwrap-test.cc index 8bb7e5c8ba06..86ca7f611196 100644 --- a/projects/rccl/test/host/rocmwrap-test.cc +++ b/projects/rccl/test/host/rocmwrap-test.cc @@ -689,8 +689,8 @@ TEST_F(RocmwrapMicrotest, Init_KernelConfigWithBothOptionsArmsTheGateAndInitsHsa EXPECT_TRUE(fopenPaths_.empty()); } -// Pinned as observed: the pfn gate is armed before the kernel check, so it stays armed on this failure. -TEST_F(RocmwrapMicrotest, Init_KernelConfigMissingEitherOptionIsASystemError) { +// The pfn gate is armed before the kernel check, so it stays armed on this failure; a fix flips this pin. +TEST_F(RocmwrapMicrotest, Init_KnownBug_KernelConfigFailureLeavesTheExportGateArmed) { for (int missing = 0; missing < 2; ++missing) { SCOPED_TRACE(missing == 0 ? "no MOVE_NOTIFY" : "no P2PDMA"); Reset(); @@ -735,8 +735,8 @@ TEST_F(RocmwrapMicrotest, Init_KallsymsFallbackMissingEitherSymbolIsASystemError } } -// Pinned as observed: an unreadable /proc/kallsyms is only logged, and init still succeeds. -TEST_F(RocmwrapMicrotest, Init_UnreadableKallsymsStillSucceeds) { +// An unreadable /proc/kallsyms is only logged, and init still succeeds; a fix flips this pin. +TEST_F(RocmwrapMicrotest, Init_KnownBug_UnreadableKallsymsStillSucceeds) { initOnceFunc(); EXPECT_TRUE(Logged(NCCL_LOG_INFO, "Could not open /proc/kallsyms")); EXPECT_EQ(initResult, ncclSuccess); diff --git a/projects/rccl/test/host/roctx-nvml-stub-test.cc b/projects/rccl/test/host/roctx-nvml-stub-test.cc index 9835b0756f74..657ed622be6d 100644 --- a/projects/rccl/test/host/roctx-nvml-stub-test.cc +++ b/projects/rccl/test/host/roctx-nvml-stub-test.cc @@ -6,7 +6,7 @@ // Host-only microtests for src/misc/roctx.cc and src/misc/nvmlwrap_stub.cc. The -roctx variant // builds roctx.cc's ROCTX_ENABLE arm; the range tests hold in either. stringify() inputs stay -// under MAX_MESSAGE_LENGTH: past it, it overflows its buffer (AICOMRCCL-2914). +// under MAX_MESSAGE_LENGTH: past it, it overflows its buffer. #include @@ -378,8 +378,8 @@ TEST_F(RoctxNvmlRoctxMicrotest, Extract_MappingToUnknownRoctxTypeTakesTheDefault roctxFree(&info); } -// Pinned as observed: an unmapped nvtx type is inserted by operator[] as INT and decoded as one. -TEST_F(RoctxNvmlRoctxMicrotest, Extract_UnmappedNvtxTypeIsTreatedAsIntAndInsertedIntoTheMap) { +// operator[] inserts an unmapped nvtx type as INT, so it is decoded as one; a fix flips this pin. +TEST_F(RoctxNvmlRoctxMicrotest, Extract_KnownBug_UnmappedNvtxTypeIsDecodedAsInt) { const uint64_t unmapped = NVTX_PAYLOAD_ENTRY_TYPE_INT + 1; ASSERT_EQ(nvtxToRoctx.count(unmapped), 0u); int value = 42;