diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index af633ee8d3c0..e329a0ce44f4 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -444,6 +444,32 @@ function(rccl_define_micro_source_lists) ../common/main_altrsmi.cpp PARENT_SCOPE) + # 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 + PARENT_SCOPE) + + # 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 + ${RCCL_HOST_HIPIFY_DIR}/src/misc/utils.cc # real busIdToInt64 + ../common/main_altrsmi.cpp + PARENT_SCOPE) + + # 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 + ${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) + # src/diagnostics/p2p.cc; fakes transport/p2p.cc, which rccl-UnitTestsMicro compiles. set(TEST_MICRO_DIAGNOSTICS_SOURCE_FILES diag-p2p-test.cc @@ -824,6 +850,8 @@ 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) + # 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. if(EXISTS "${RCCL_HIPIFY_DIR}/gensrc/symmetric/sym_kernels_host.cc") @@ -860,6 +888,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) --- @@ -941,7 +970,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 @@ -1203,7 +1232,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 @@ -1243,6 +1272,46 @@ rccl_add_micro_binary(rccl-UnitTestsMicroRomeModels ROME_MODELS_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/graph/rome_models.cc" CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_rome_models.yaml") +# --------------------------------------------------------------------------- +# rccl-UnitTestsMicroRoctxNvml +# --------------------------------------------------------------------------- +set(_roctx_nvml_common_args + 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_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. +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-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-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-UnitTestsMicroGraph is its own binary because the others link fakes for the graph symbols these TUs define. # --------------------------------------------------------------------------- diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index fc128b82206e..7b8309fdce7f 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -350,6 +350,21 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: memory stack. The suite drives the public lifecycle -- `ncclRmaProxyCreateContext` / `ncclRmaProxyDestroyContext` -- so production owns the context and its cleanup. See `test_categories_micro_rma_proxy.yaml`. +- **`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. See + `test_categories_micro_roctx_nvml.yaml`. +- **`rccl-UnitTestsMicroRocmwrap`** — `src/misc/rocmwrap.cc` (via + `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-UnitTestsMicroAmdsmiWrap`** — `src/misc/amdsmi_wrap.cc` (via + `AMDSMI_WRAP_CC_PATH`, from `amdsmi-wrap-test.cc`); suite + `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` @@ -882,15 +897,15 @@ can be configured **directly** to build every host binary — `rccl-HostUnitTest 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 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. `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. @@ -927,6 +942,10 @@ cmake --build build -j"$(nproc)" ./build/rccl-UnitTestsMicroRmaProxy # src/rma/rma_proxy.cc tests ./build/rccl-UnitTestsMicroDda # src/algorithms/dda/ tests ./build/rccl-HostUnitTests +./build/rccl-UnitTestsMicroRoctxNvml # misc/roctx.cc + misc/nvmlwrap_stub.cc tests +./build/rccl-UnitTestsMicroRoctxNvml-roctx # same, ROCTX_ENABLE arm +./build/rccl-UnitTestsMicroRocmwrap # src/misc/rocmwrap.cc cuMem + DMA-BUF init +./build/rccl-UnitTestsMicroAmdsmiWrap # src/misc/amdsmi_wrap.cc ``` Disable coverage instrumentation for the standalone host-only test binaries 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..e90424df554a --- /dev/null +++ b/projects/rccl/test/host/amdsmi-wrap-test.cc @@ -0,0 +1,1375 @@ +/************************************************************************* + * 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 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 + +#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; +} // 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; +} + +#define access AmdsmiTestAccess +#define dlopen AmdsmiTestDlopen +#define dlsym AmdsmiTestDlsym +#include AMDSMI_WRAP_CC_PATH +#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); + +// 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) { + 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 procsCalls = 0; + int typeCalls = 0; + int enumCalls = 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) { + ++g_smi.procsCalls; + 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) { + ++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; + 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; } + +// 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() { + for (const auto& sym : SmiSymbols()) *sym.slot = sym.fake; +} + +void ClearSmiTable() { + for (const auto& sym : SmiSymbols()) *sym.slot = 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_legacyTelemName = "LEGACY_NAME"; + 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; }; + 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 + +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; + // 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 +// 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; 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; + 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]; the source comment says +// the domain is [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); +} + +// `(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; + 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); + // 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) { + 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(); + 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 + 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); + 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, 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; + 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(); + 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); +} + +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")); + Reset(); // the boundary itself fits + 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 (!kTypedFabricLayout) { + 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 (!kTypedFabricLayout) 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 (!kTypedFabricLayout) 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 && !kTypedFabricLayout) 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); + // 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); + + 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/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/rocmwrap-test.cc b/projects/rccl/test/host/rocmwrap-test.cc new file mode 100644 index 000000000000..86ca7f611196 --- /dev/null +++ b/projects/rccl/test/host/rocmwrap-test.cc @@ -0,0 +1,779 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// 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 + +#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 = [this](const char* path, const char*) -> FILE* { + fopenPaths_.push_back(path); + 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 >= ROCM_VER_7_12_60540 + 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{}; + std::snprintf(prop->gcnArchName, sizeof(prop->gcnArchName), "gfx942:sramecc+:xnack-"); + 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; + // 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")); + 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); + // net.cc and coll_net.cc gate on the pointer itself, not on initResult. + EXPECT_EQ(pfn_hsa_amd_portable_export_dmabuf, nullptr); + } +} + +// 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()); +} + +// 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(); + 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); + } +} + +// 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); + 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/roctx-nvml-stub-test.cc b/projects/rccl/test/host/roctx-nvml-stub-test.cc new file mode 100644 index 000000000000..657ed622be6d --- /dev/null +++ b/projects/rccl/test/host/roctx-nvml-stub-test.cc @@ -0,0 +1,583 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// 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. + +#include + +#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's own bodies only, so roctxAlloc's failure paths are reachable. +#include +#include "roctx.h" +namespace { +std::function g_roctxMalloc; +} // namespace + +// With ROCTX_ENABLE, record roctracer's push/pop instead of linking it. +#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 +#undef malloc + +#include NVMLWRAP_STUB_CC_PATH + +namespace { + +class RoctxNvmlRoctxMicrotest : public ::testing::Test { +protected: + void SetUp() override { + ResetRecorders(); + savedNvtxToRoctx_ = nvtxToRoctx; + } + void TearDown() override { + g_loadParam = [](const char*, int64_t deftVal) { return deftVal; }; + g_roctxMalloc = nullptr; + // Restored here so a failed assertion can't leak entries into the next test. + nvtxToRoctx = savedNvtxToRoctx_; + ResetRecorders(); + } + + static void ResetRecorders() { + g_logged.clear(); +#ifdef ROCTX_ENABLE + g_roctxPushed.clear(); + g_roctxPops = 0; +#endif + } + + // 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"; + 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 + } + + // 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"; + 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; + 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; + } + + // 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; + std::memset(&e, 0, sizeof(e)); + e.type = type; + e.name = name; + e.offset = offset; + return e; + } + +private: + std::map savedNvtxToRoctx_; +}; + +struct MixedPayload { + int i; + size_t s; + ncclDevRedOp_t op; + ncclDataType_t dt; +}; + +} // namespace + +// Lookup tables + +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); + 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(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; + + 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; +} + +// 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"); + 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"); +} + +// roctxAlloc / roctxFree + +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 (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(RoctxNvmlRoctxMicrotest, 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: entries are freed and nulled. +TEST_F(RoctxNvmlRoctxMicrotest, 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(RoctxNvmlRoctxMicrotest, 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(RoctxNvmlRoctxMicrotest, 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(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)), + 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(RoctxNvmlRoctxMicrotest, Stringify_SingleEntryHasNoSeparator) { + EXPECT_EQ(Stringify("Bcast", {IntEntry("root", 7)}), "{Bcast: root: 7}"); +} + +TEST_F(RoctxNvmlRoctxMicrotest, Stringify_ZeroEntriesIsJustTheId) { + EXPECT_EQ(Stringify("Empty", {}), "{Empty: }"); +} + +// 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, {}}; + 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(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(RoctxNvmlRoctxMicrotest, 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 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; + info.payloadEntries = &entry; + info.id = "X"; + info.numEntries = 1; + stringify(&info); + EXPECT_EQ(info.message, nullptr); +} + +// 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); + 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); + EXPECT_EQ(info.payloadEntries[0].type, bogus); + EXPECT_STREQ(info.message, "{D: w: unknown roctx payload type}"); + roctxFree(&info); +} + +// 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; + 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); // TearDown restores the table + roctxFree(&info); +} + +// roctx_scoped_range_in + +TEST_F(RoctxNvmlRoctxMicrotest, 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(RoctxNvmlRoctxMicrotest, ScopedRange_LoggingOff_DoesNothing) { + SetLogRoctx(0); + { + roctx_scoped_range_in range("msg"); + roctx_scoped_range_in defaulted; + } + EXPECT_EQ(RangeEvents(), 0u); + EXPECT_TRUE(g_logged.empty()); +} + +// One range: one enter event and one exit event. +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_LoggingOn_OneEventOnEnterAndOneOnExit) { + SetLogRoctx(1); + { + roctx_scoped_range_in range("msg"); + EXPECT_EQ(RangeEvents(), 1u); + ExpectPushedMessage("msg"); + } + ExpectOneRangeEnteredAndExited(); +} + +TEST_F(RoctxNvmlRoctxMicrotest, ScopedRange_DefaultConstructorDelegatesToTheMessageConstructor) { + SetLogRoctx(1); + { + roctx_scoped_range_in range; + ExpectPushedMessage(""); + } + ExpectOneRangeEnteredAndExited(); +} + +// 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(); + 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}"); + } + ExpectOneRangeEnteredAndExited(); +} + +// 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); + 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(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()); +} + +// nvmlwrap_stub.cc: no NVML on ROCm, so every query reports "no device". + +namespace { +nvmlDevice_t FakeDevice() { return reinterpret_cast(0x1234); } +} // namespace + +TEST(RoctxNvmlStubMicrotest, LifecycleCallsSucceed) { + EXPECT_EQ(ncclNvmlSymbols(), ncclSuccess); + EXPECT_EQ(ncclNvmlInit(), ncclSuccess); + EXPECT_EQ(ncclNvmlShutdown(), ncclSuccess); +} + +TEST(RoctxNvmlStubMicrotest, 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(RoctxNvmlStubMicrotest, 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(RoctxNvmlStubMicrotest, 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. +TEST(RoctxNvmlStubMicrotest, 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(RoctxNvmlStubMicrotest, 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(RoctxNvmlStubMicrotest, 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 16aa52c13e30..6a69aefce8d2 100755 --- a/projects/rccl/test/host/run_host_tests.sh +++ b/projects/rccl/test/host/run_host_tests.sh @@ -170,6 +170,11 @@ do_host_tests() { "rccl-UnitTestsMicroClientSupport:$SCRIPT_DIR/host_tests_micro_client_support.xml" "rccl-UnitTestsMicroGinHost:$SCRIPT_DIR/host_tests_micro_gin_host.xml" "rccl-UnitTestsMicroDdaIpc:$SCRIPT_DIR/host_tests_micro_dda_ipc.xml" + "rccl-UnitTestsMicroRoctxNvml:$SCRIPT_DIR/host_tests_micro_roctx_nvml.xml" + # Standalone-only: roctx.cc's ROCTX_ENABLE arm. + "rccl-UnitTestsMicroRoctxNvml-roctx:$SCRIPT_DIR/host_tests_micro_roctx_nvml_roctx.xml" + "rccl-UnitTestsMicroRocmwrap:$SCRIPT_DIR/host_tests_micro_rocmwrap.xml" + "rccl-UnitTestsMicroAmdsmiWrap:$SCRIPT_DIR/host_tests_micro_amdsmi_wrap.xml" "rccl-UnitTestsMicroRomeModels:$SCRIPT_DIR/host_tests_micro_rome_models.xml" "rccl-UnitTestsMicroDiagnostics:$SCRIPT_DIR/host_tests_micro_diagnostics.xml" "rccl-UnitTestsMicroGraph:$SCRIPT_DIR/host_tests_micro_graph.xml" 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/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/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 c10ff495eb51..1abecfa3b2d7 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json +++ b/projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json @@ -287,6 +287,39 @@ {"name": "DdaIpcMicrotests", "description": "Host-only dda_all_reduce_ipc.cu eligibility, datatype dispatch and flat/tree selection microtests (no GPU, no HIP/RCCL runtime)", "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": "*"} + ] + }, + "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": "*"} + ] + }, + "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": "*"} + ] + }, "unit_tests_micro_rome_models": { "is_gtest": true, "binary": "host/rccl-UnitTestsMicroRomeModels", @@ -383,6 +416,24 @@ "config": "unit_tests_micro_dda_ipc", "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 + }, + { + "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 + }, + { + "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 + }, { "name": "RCCL Microtests - rome_models.cc", "description": "rccl-UnitTestsMicroRomeModels (host-only rome_models.cc microtests, no GPU/HIP/RCCL runtime)", 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 6f0622289d64..65427fdd1ee1 100644 --- a/projects/rccl/tools/scripts/test_runner/lib/test_executor.py +++ b/projects/rccl/tools/scripts/test_runner/lib/test_executor.py @@ -3288,6 +3288,38 @@ def generate_coverage_report(self): # llvm-cov to attribute that coverage. host_test_dir = os.path.join(test_dir, "host") for binary in ["rccl-UnitTestsMicro", "rccl-UnitTestsMicroInit", + "rccl-UnitTestsMicroAlloc", + "rccl-UnitTestsMicroAllocator", + "rccl-UnitTestsMicroAmdsmiWrap", + "rccl-UnitTestsMicroApiTrace", + "rccl-UnitTestsMicroArgcheck", + "rccl-UnitTestsMicroBootstrap", + "rccl-UnitTestsMicroCftDevRuntime", + "rccl-UnitTestsMicroCollTrace", + "rccl-UnitTestsMicroCollectives", + "rccl-UnitTestsMicroCore", + "rccl-UnitTestsMicroCpuset", + "rccl-UnitTestsMicroCrypt", + "rccl-UnitTestsMicroDdaColl", + "rccl-UnitTestsMicroDebug", + "rccl-UnitTestsMicroGinHostProxy", + "rccl-UnitTestsMicroLinux", + "rccl-UnitTestsMicroMemManager", + "rccl-UnitTestsMicroMgmtTaskEnq", + "rccl-UnitTestsMicroMnnvl", + "rccl-UnitTestsMicroNcclDevice", + "rccl-UnitTestsMicroOnerank", + "rccl-UnitTestsMicroParam", + "rccl-UnitTestsMicroPlugin", + "rccl-UnitTestsMicroProfiler", + "rccl-UnitTestsMicroProxy", + "rccl-UnitTestsMicroRecorder", + "rccl-UnitTestsMicroRocmwrap", + "rccl-UnitTestsMicroRoctxNvml", + "rccl-UnitTestsMicroRomeModels", + "rccl-UnitTestsMicroStrongstream", + "rccl-UnitTestsMicroTopo", + "rccl-UnitTestsMicroUtils", "rccl-UnitTestsMicroInit-uncached", "rccl-UnitTestsMicroInit-faultinj", "rccl-UnitTestsMicroEnqueue",