pvallem · Oct 17, 2026 · Oct 8, 2026 · Oct 15, 2026 · Oct 15, 2026 · Oct 16, 2026
@@ -392,6 +392,28 @@ function(rccl_define_micro_source_lists)
     ../common/main_altrsmi.cpp
     PARENT_SCOPE)
 
+  # rccl-UnitTestsMicroGraph: the real src/graph topology TUs, which every other micro binary fakes.
+  set(TEST_MICRO_GRAPH_SOURCE_FILES
+    graph-dev-test.cc
+    fakes/amdsmi_fakes.cc
+    fakes/env_fakes.cc
+    fakes/hip_fakes.cc
+    fakes/nccl_fakes.cc
+    fakes/os_fakes.cc
+    fakes/rccl_wrap_fakes.cc
+    fakes/transport_stubs.cc
+    fakes/tuning_fakes.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/graph/paths.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/graph/rome_models.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/graph/search.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/graph/topo.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/graph/xml.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/archinfo.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/nvmlwrap_stub.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/utils.cc
+    ../common/main_altrsmi.cpp
+    PARENT_SCOPE)
+
   # rccl-UnitTestsMicroGinHost -- src/gin/gin_host.cc (GIN_PROXY_NTHREADS, proxy affinity). Own
   # binary: gin-plugin-init-test.cc already defines ncclParamGinEnable in
   # rccl-UnitTestsMicro, and gin_fakes.cc supplies ncclGinQueryLastError to Init.
@@ -1185,6 +1207,26 @@ 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-UnitTestsMicroGraph is its own binary because the others link fakes for the graph symbols these TUs define.
+# ---------------------------------------------------------------------------
+set_source_files_properties(
+  ${RCCL_HOST_HIPIFY_DIR}/src/graph/paths.cc
+  ${RCCL_HOST_HIPIFY_DIR}/src/graph/rome_models.cc
+  ${RCCL_HOST_HIPIFY_DIR}/src/graph/search.cc
+  ${RCCL_HOST_HIPIFY_DIR}/src/graph/topo.cc
+  ${RCCL_HOST_HIPIFY_DIR}/src/graph/xml.cc
+  ${RCCL_HOST_HIPIFY_DIR}/src/misc/nvmlwrap_stub.cc
+  PROPERTIES GENERATED TRUE)
+rccl_add_micro_binary(rccl-UnitTestsMicroGraph
+  SOURCES ${TEST_MICRO_GRAPH_SOURCE_FILES}
+  EXTRA_DEFS
+    ${_MICRO_MODE_DEFS}
+    ${_MICRO_AMDSMI_DEF}
+    RCCL_NCCL_FAKES_OMIT_GRAPH
+    RCCL_TRANSPORT_STUBS_OMIT_ncclPxnDisable
+  CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_graph.yaml")
+
 # ---------------------------------------------------------------------------
 # rccl-UnitTestsMicroDiagnostics
 # ---------------------------------------------------------------------------

@@ -342,6 +342,10 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary:
   (via `DIAG_IB_WRITE_BW_CC_PATH`, suite `DiagIbWriteBwMicrotest.*`). Its own binary: it fakes the
   `transport/p2p.cc` shareable-buffer entry points that `rccl-UnitTestsMicro`
   compiles for real. See `test_categories_micro_diagnostics.yaml`.
+- **`rccl-UnitTestsMicroGraph`**: the real `src/graph/{xml,topo,paths,search,rome_models}.cc`,
+  compiled as sources; suite `TopoGpuDevMicrotest.*` in `graph-dev-test.cc`. Its own binary:
+  every other target links fakes for the graph symbols these files define. See
+  `test_categories_micro_graph.yaml`.
 
 Everything below (seams, fakes, coverage) applies to both; the concrete examples
 use `p2p.cc`.
@@ -857,6 +861,7 @@ can be configured **directly** to build every host binary — `rccl-HostUnitTest
 `rccl-UnitTestsMicro`, `rccl-UnitTestsMicroWarpSpeed`,
 `rccl-UnitTestsMicroInit[-uncached|-faultinj]`, `rccl-UnitTestsMicroEnqueue[-devlinker]`,
 `rccl-UnitTestsMicroGinHost`, `rccl-UnitTestsMicroCe`, `rccl-UnitTestsMicroDiagnostics`,
+`rccl-UnitTestsMicroGraph`,
 `rccl-UnitTestsMicroSymKernels`, `rccl-UnitTestsMicroTaskPrep` and
 `rccl-UnitTestsMicroRmaProxy` — **without configuring/building all of
 librccl**. It compiles just the tests + fakes + the hipified unit-under-test
@@ -898,6 +903,7 @@ cmake --build build -j"$(nproc)"
 ./build/rccl-UnitTestsMicroTaskPrep           # src/enqueue/task_prep/ + task_sched/ tests
 ./build/rccl-UnitTestsMicroConnect            # src/graph/connect.cc Preset/Postset tests
 ./build/rccl-UnitTestsMicroDiagnostics        # src/diagnostics/{p2p,ib_write_bw}.cc tests
+./build/rccl-UnitTestsMicroGraph              # src/graph/{xml,topo,paths,search,rome_models}.cc tests
 ./build/rccl-UnitTestsMicroGinHost            # src/gin/gin_host.cc GIN_PROXY_NTHREADS
 ./build/rccl-UnitTestsMicroDiagnostics        # src/diagnostics/p2p.cc tests
 ./build/rccl-UnitTestsMicroCe                 # src/ce_coll.cc tests

@@ -7,6 +7,15 @@
 // Implementation of the src/misc/amdsmi_wrap.cc fakes. See amdsmi_fakes.h.
 
 #include "amdsmi_fakes.h"
+#include "signature-drift.h"
+
+ASSERT_HOOK_MATCHES_PROD(g_amdSmiGetDeviceIndexByPciBusId, amd_smi_getDeviceIndexByPciBusId);
+ASSERT_HOOK_MATCHES_PROD(g_amdSmiGetFabricDeviceInfo,      amd_smi_getFabricDeviceInfo);
+ASSERT_HOOK_MATCHES_PROD(g_amdSmiGetNumDevice,             amd_smi_getNumDevice);
+ASSERT_HOOK_MATCHES_PROD(g_amdSmiGetDevicePciBusIdString,  amd_smi_getDevicePciBusIdString);
+ASSERT_HOOK_MATCHES_PROD(g_amdSmiGetLinkInfo,              amd_smi_getLinkInfo);
+ASSERT_HOOK_MATCHES_PROD(g_amdSmiFabricTelemetryAcquire,   amd_smi_fabricTelemetryAcquire);
+ASSERT_HOOK_MATCHES_PROD(g_amdSmiFabricTelemetryRelease,   amd_smi_fabricTelemetryRelease);
 
 ncclResult_t DefaultAmdSmiGetDeviceIndexByPciBusId(const char*, uint32_t* deviceIndex) {
   if (deviceIndex) *deviceIndex = static_cast<uint32_t>(-1);  // -1 -> skip fabric block
@@ -30,6 +39,26 @@ ncclResult_t amd_smi_getFabricDeviceInfo(uint32_t deviceIndex, struct amdsmiFabr
 ncclResult_t g_amdSmiInitResult = ncclSuccess;
 ncclResult_t amd_smi_init() { return g_amdSmiInitResult; }
 
+static ncclResult_t DefaultAmdSmiGetNumDevice(uint32_t* numDevs) {
+  *numDevs = 0;
+  return ncclSuccess;
+}
+std::function<ncclResult_t(uint32_t*)> g_amdSmiGetNumDevice = DefaultAmdSmiGetNumDevice;
+ncclResult_t amd_smi_getNumDevice(uint32_t* numDevs) { return g_amdSmiGetNumDevice(numDevs); }
+
+static ncclResult_t DefaultAmdSmiGetDevicePciBusIdString(uint32_t, char*, size_t) { return ncclInternalError; }
+std::function<ncclResult_t(uint32_t, char*, size_t)> g_amdSmiGetDevicePciBusIdString =
+    DefaultAmdSmiGetDevicePciBusIdString;
+ncclResult_t amd_smi_getDevicePciBusIdString(uint32_t deviceIndex, char* pciBusId, size_t len) {
+  return g_amdSmiGetDevicePciBusIdString(deviceIndex, pciBusId, len);
+}
+
+static ncclResult_t DefaultAmdSmiGetLinkInfo(int, int, amdsmi_link_type_t*, int*, int*) { return ncclInternalError; }
+std::function<ncclResult_t(int, int, amdsmi_link_type_t*, int*, int*)> g_amdSmiGetLinkInfo = DefaultAmdSmiGetLinkInfo;
+ncclResult_t amd_smi_getLinkInfo(int srcDev, int dstDev, amdsmi_link_type_t* type, int* hops, int* count) {
+  return g_amdSmiGetLinkInfo(srcDev, dstDev, type, hops, count);
+}
+
 ncclResult_t DefaultAmdSmiFabricTelemetryAcquire(uint32_t, uint64_t, int, bool* acquired) {
   if (acquired) *acquired = false;  // telemetry is opt-in and off by default
   return ncclSuccess;
@@ -50,6 +79,9 @@ void ResetAmdSmiFakes() {
   g_amdSmiGetDeviceIndexByPciBusId = DefaultAmdSmiGetDeviceIndexByPciBusId;
   g_amdSmiGetFabricDeviceInfo = DefaultAmdSmiGetFabricDeviceInfo;
   g_amdSmiInitResult = ncclSuccess;
+  g_amdSmiGetNumDevice = DefaultAmdSmiGetNumDevice;
+  g_amdSmiGetDevicePciBusIdString = DefaultAmdSmiGetDevicePciBusIdString;
+  g_amdSmiGetLinkInfo = DefaultAmdSmiGetLinkInfo;
   g_amdSmiFabricTelemetryAcquire = DefaultAmdSmiFabricTelemetryAcquire;
   g_amdSmiFabricTelemetryRelease = DefaultAmdSmiFabricTelemetryRelease;
 }
@@ -13,10 +13,9 @@
 #include <cstdint>
 #include <functional>
 
+#include "amdsmi_wrap.h"
 #include "nccl.h"
 
-struct amdsmiFabricDeviceInfo;
-
 // fillInfo's UALoE/MNNVL probe. The default answers -1, i.e. what a host with no fabric device reports.
 ncclResult_t DefaultAmdSmiGetDeviceIndexByPciBusId(const char* busId, uint32_t* deviceIndex);
 extern std::function<ncclResult_t(const char*, uint32_t*)> g_amdSmiGetDeviceIndexByPciBusId;
@@ -27,6 +26,11 @@ extern std::function<ncclResult_t(uint32_t, struct amdsmiFabricDeviceInfo*)> g_a
 
 extern ncclResult_t g_amdSmiInitResult;
 
+// XGMI discovery in ncclTopoGetXmlFromGpu. The defaults report an empty machine and fail every per-device query.
+extern std::function<ncclResult_t(uint32_t*)> g_amdSmiGetNumDevice;
+extern std::function<ncclResult_t(uint32_t, char*, size_t)> g_amdSmiGetDevicePciBusIdString;
+extern std::function<ncclResult_t(int, int, amdsmi_link_type_t*, int*, int*)> g_amdSmiGetLinkInfo;
+
 // commAlloc's fabric telemetry opt-in. The default declines, i.e. what every host
 // reports unless RCCL_FABRIC_TELEMETRY_ENABLE is set, which leaves commFree with
 // nothing to release.

@@ -157,15 +157,6 @@ ncclResult_t ncclCommWindowDeregister(ncclComm_t comm, ncclWindow_t win) {
 // Group state machine.
 // ---------------------------------------------------------------------------
 
-// ---------------------------------------------------------------------------
-// Param loader.
-// ---------------------------------------------------------------------------
-int64_t ncclLoadParam(char const*, int64_t deftVal, int64_t, int64_t* cache, int8_t* noCache) {
-  if (cache) *cache = deftVal;
-  if (noCache) *noCache = 0;
-  return deftVal;
-}
-
 // Emitted by init.cc, which this binary does not link, so the redirected
 // NCCL_PARAM in dev_runtime.cc does not cover it. Routed through g_loadParam
 // like the params dev_runtime.cc owns, so a test can drive the NCCL_WIN_ENABLE=0

@@ -52,6 +52,7 @@ ASSERT_HOOK_MATCHES_PROD(g_hipMemMap,                     hipMemMap);
 ASSERT_HOOK_MATCHES_PROD(g_hipMemSetAccess,               hipMemSetAccess);
 ASSERT_HOOK_MATCHES_PROD(g_hipIpcOpenMemHandle,           hipIpcOpenMemHandle);
 ASSERT_HOOK_MATCHES_PROD(g_hipDeviceGetPCIBusId,          hipDeviceGetPCIBusId);
+ASSERT_HOOK_MATCHES_PROD(g_hipDeviceGetByPCIBusId,        hipDeviceGetByPCIBusId);
 ASSERT_HOOK_MATCHES_PROD(g_hipEventRecord,                hipEventRecord);
 ASSERT_HOOK_MATCHES_PROD(g_hipStreamBatchMemOp,           hipStreamBatchMemOp);
 ASSERT_HOOK_MATCHES_PROD(g_hipStreamWriteValue64,         hipStreamWriteValue64);
@@ -244,6 +245,7 @@ std::function<hipError_t(int*)> g_hipGetDeviceCount = DefaultHipGetDeviceCount;
 static hipError_t DefaultHipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device);
 static hipError_t DefaultHipDeviceSetLimit(hipLimit_t limit, size_t value);
 static hipError_t DefaultHipDeviceGetPCIBusId(char* pciBusId, int len, int device);
+static hipError_t DefaultHipDeviceGetByPCIBusId(int* device, const char* pciBusId);
 static hipError_t DefaultHipEventRecord(hipEvent_t event, hipStream_t stream);
 
 static hipError_t DefaultHipDeviceCanAccessPeer(int* canAccessPeer, int, int)
@@ -631,6 +633,7 @@ void ResetHipFakes()
     g_hipDeviceGetAttributeResult   = hipErrorInvalidValue;
     g_hipDeviceGetPCIBusIdResult    = hipErrorInvalidValue;
     g_hipDeviceGetPCIBusId          = DefaultHipDeviceGetPCIBusId;
+    g_hipDeviceGetByPCIBusId        = DefaultHipDeviceGetByPCIBusId;
     g_hipEventCreateResult          = hipErrorInvalidValue;
     g_hipMemPoolResult              = hipErrorInvalidValue;
     g_hipStreamCreateResult         = hipErrorInvalidValue;
@@ -776,6 +779,13 @@ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device)
     return g_hipDeviceGetPCIBusId(pciBusId, len, device);
 }
 
+static hipError_t DefaultHipDeviceGetByPCIBusId(int*, const char*) { return hipErrorInvalidValue; }
+std::function<hipError_t(int*, const char*)> g_hipDeviceGetByPCIBusId = DefaultHipDeviceGetByPCIBusId;
+hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId)
+{
+    return g_hipDeviceGetByPCIBusId(device, pciBusId);
+}
+
 hipError_t hipEventCreate(hipEvent_t* event)
 {
     if (event) {

@@ -109,6 +109,8 @@ extern hipError_t g_hipDeviceGetPCIBusIdResult;
 // device index (drives busIdToCudaDev's per-device resolution).
 extern std::function<hipError_t(char* /*pciBusId*/, int /*len*/, int /*device*/)>
     g_hipDeviceGetPCIBusId;
+// Backs hipDeviceGetByPCIBusId. The default fails, as for a BDF this process cannot see.
+extern std::function<hipError_t(int* /*device*/, const char* /*pciBusId*/)> g_hipDeviceGetByPCIBusId;
 extern hipError_t g_hipEventCreateResult;
 extern hipError_t g_hipMemPoolResult;
 extern hipError_t g_hipStreamCreateResult;

@@ -87,14 +87,15 @@ void ncclDebugLog(ncclDebugLogLevel /*level*/,
     std::fputc('\n', stderr);
 }
 
-void ncclLoadParam(char const* /*env*/,
-                   int64_t     /*deftVal*/,
-                   int64_t     /*uninitialized*/,
-                   int64_t*    /*cache*/)
-{
-    // No-op: leaves cache untouched so NCCL_PARAM callers in non-shimmed
-    // TUs see the default. The NCCL_PARAM bodies that p2p-test.cc
-    // redirects through g_loadParam bypass this entirely.
+// Every param sits at its default. The NCCL_PARAM bodies that p2p-test.cc redirects through g_loadParam bypass this.
+int64_t ncclLoadParam(char const*, int64_t deftVal, int64_t, int64_t* cache, int8_t* noCache) {
+  if (cache) {
+    *cache = deftVal;
+  }
+  if (noCache) {
+    *noCache = 0;
+  }
+  return deftVal;
 }
 
 // Default returns deftVal verbatim -- preserves the pre-hook contract that
@@ -260,6 +261,8 @@ static ncclResult_t DefaultTopoCheckP2p(int /*rank1*/, int /*rank2*/, int* p2p,
 std::function<ncclResult_t(int, int, int*, int*, int*, int*, int*)>
     g_ncclTopoCheckP2p = DefaultTopoCheckP2p;
 
+// rccl-UnitTestsMicroGraph links the real src/graph/paths.cc and search.cc, which define these.
+#ifndef RCCL_NCCL_FAKES_OMIT_GRAPH
 ncclResult_t ncclTopoCheckP2p(struct ncclComm*       /*comm*/,
                               struct ncclTopoSystem* /*system*/,
                               int                    rank1,
@@ -276,6 +279,7 @@ ncclResult_t ncclTopoCheckP2p(struct ncclComm*       /*comm*/,
     if (isCrossClique) *isCrossClique = 0;
     return g_ncclTopoCheckP2p(rank1, rank2, p2p, read, intermediateRank, cudaP2p, isCrossClique);
 }
+#endif
 
 // --- Controllable seam: ncclTopoCheckNet ----------------------------------
 // Default preserves the old stub: report NET is not better (net = 0).
@@ -287,13 +291,15 @@ static ncclResult_t DefaultTopoCheckNet(int /*rank1*/, int /*rank2*/, int* net)
 std::function<ncclResult_t(int, int, int*)>
     g_ncclTopoCheckNet = DefaultTopoCheckNet;
 
+#ifndef RCCL_NCCL_FAKES_OMIT_GRAPH
 ncclResult_t ncclTopoCheckNet(struct ncclTopoSystem* /*system*/,
                               int                    rank1,
                               int                    rank2,
                               int*                   net)
 {
     return g_ncclTopoCheckNet(rank1, rank2, net);
 }
+#endif
 
 // --- Controllable seam: ncclCommGraphRegister ---------------------------
 // Default preserves the old stub: fail with a null handle. Graph-register
@@ -418,17 +424,19 @@ ncclResult_t DefaultTopoGetLinkType(int, int, bool* isXGMI, int)
 std::function<ncclResult_t(int, int, bool*, int)> g_ncclTopoGetLinkType = DefaultTopoGetLinkType;
 int g_ncclTopoGetLinkTypeCalls = 0;
 
+#ifndef RCCL_NCCL_FAKES_OMIT_GRAPH
 ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* /*system*/,
-                                 int                    cudaDev1,
-                                 int                    cudaDev2,
+                                 int                    dev1,
+                                 int                    dev2,
                                  bool*                  isXGMI,
                                  int                    maxInter,
                                  int                    /*nInter*/,
                                  int*                   /*inter*/)
 {
     g_ncclTopoGetLinkTypeCalls++;
-    return g_ncclTopoGetLinkType(cudaDev1, cudaDev2, isXGMI, maxInter);
+    return g_ncclTopoGetLinkType(dev1, dev2, isXGMI, maxInter);
 }
+#endif
 
 // ---------------------------------------------------------------------------
 // Memory-manager tracking stubs.

@@ -119,7 +119,7 @@ extern std::function<int64_t(const char* /*env*/, int64_t /*deftVal*/)>
 
 // The two device ids and the isXGMI out-parameter; the topo system and the inter-GPU knobs are dropped.
 // maxInter is carried because init.cc:1834 passes an explicit 1 where graph.h:90 defaults to MAX_XGMI_INTER_GPUS.
-extern std::function<ncclResult_t(int /*cudaDev1*/, int /*cudaDev2*/, bool* /*isXGMI*/, int /*maxInter*/)>
+extern std::function<ncclResult_t(int /*dev1*/, int /*dev2*/, bool* /*isXGMI*/, int /*maxInter*/)>
     g_ncclTopoGetLinkType;
 extern int g_ncclTopoGetLinkTypeCalls;
 

@@ -111,6 +111,27 @@ ncclResult_t ncclOsGetPciDeviceComputePartitionByBusId(const char* busId, char*
   return ncclSuccess;
 }
 
+void ncclOsCpuZero(ncclAffinity& affinity) { CPU_ZERO(&affinity); }
+void ncclOsCpuSet(ncclAffinity& affinity, int cpu) { CPU_SET(cpu, &affinity); }
+
+// Topology discovery's sysfs reads. With no PCI path, ncclTopoGetXmlFromSys keeps whatever attributes a test pre-built.
+ncclResult_t ncclOsGetPciPath(const char*, char** path) {
+  *path = nullptr;
+  return ncclSystemError;
+}
+ncclResult_t ncclOsGetBcmLinks(const char*, int* nlinks, char** peers) {
+  *nlinks = 0;
+  *peers = nullptr;
+  return ncclSuccess;
+}
+ncclResult_t ncclOsGetNumaNodeAffinity(unsigned int, char* affinityStr, size_t maxLen, int* cpuOffset) {
+  *cpuOffset = 0;
+  if (maxLen > 0) {
+    affinityStr[0] = '\0';
+  }
+  return ncclSuccess;
+}
+
 void ResetOsFakes() {
   g_ncclOsCpuCountValue = 0;
   g_ncclOsCpuCountCalls = 0;

@@ -234,6 +234,7 @@ ncclResult_t rcclCommSetP2pShiftSize(struct ncclComm*) {
 int64_t rcclParamHierarchicalLazyInit() {
   FailLoudUnfaked("rccl_wrap_fakes", "rcclParamHierarchicalLazyInit");
 }
+void rcclSetPxn(struct ncclComm*, int&) { FailLoudUnfaked("rccl_wrap_fakes", "rcclSetPxn"); }
 int64_t g_rcclParamDirectReduceScatterThreshold = 8388608;     // rccl_wrap.cc:51 default
 int64_t rcclParamDirectReduceScatterThreshold() {              // rccl_wrap.cc:51
   return g_rcclParamDirectReduceScatterThreshold;

@@ -57,7 +57,9 @@ bool rcclUseAinic() { return g_useAinic(); }
 
 static int DefaultPxnDisable(struct ncclComm*) { return 0; }
 std::function<int(struct ncclComm*)> g_pxnDisable = DefaultPxnDisable;
+#ifndef RCCL_TRANSPORT_STUBS_OMIT_ncclPxnDisable
 int ncclPxnDisable(struct ncclComm* comm) { return g_pxnDisable(comm); }
+#endif
 
 // src/proxy.cc: comm teardown stops the proxy through this, so it needs a working default, not a fail-loud one.
 static ncclResult_t DefaultNcclProxyStop(struct ncclComm*) { return ncclSuccess; }
@@ -134,3 +136,12 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
 // a definition even though the code path that uses it is dead.
 #include "transport.h"
 struct ncclTransport netTransport = {};
+
+// src/transport.cc's table. ncclTopoComputePaths asks P2P, then SHM, whether a peer is reachable; both say yes.
+static ncclResult_t ReachableCanConnect(int* ret, struct ncclComm*, struct ncclTopoGraph*, struct ncclPeerInfo*,
+                                        struct ncclPeerInfo*) {
+  *ret = 1;
+  return ncclSuccess;
+}
+static struct ncclTransport reachableTransport = {"fake", ReachableCanConnect};
+struct ncclTransport* ncclTransports[NTRANSPORTS] = {&reachableTransport, &reachableTransport};
@@ -67,6 +67,9 @@ ncclResult_t ncclTuningCompute(struct ncclTuningInput_t* const input, struct ncc
   return g_tuningCompute(input, result);
 }
 
+// A no-op: the overrides only retune system->tuning after a Rome model matches.
+void rcclApplyTuningOverrides(struct ncclTopoSystem*) {}
+
 void ResetTuningFakes() {
   g_topoGetAlgoTime = DefaultTopoGetAlgoTime;
   g_topoGetAlgoTimeCalls = 0;