aduforest · Oct 12, 2026 · Oct 13, 2026 · Oct 13, 2026 · Oct 13, 2026 · Oct 14, 2026
@@ -293,6 +293,39 @@ function(rccl_define_micro_source_lists)
     ../common/main_altrsmi.cpp
     PARENT_SCOPE)
 
+  # rccl-UnitTestsMicroNcclDeviceCft -- src/nccl_device/core.cc + src/cft_dev_runtime.cc.
+  # Its own binary: rccl-UnitTestsMicro's fakes/dev_runtime_micro_fakes.cc, fakes/devcomm_fakes.cc,
+  # and gin-host-test.cc between them fake 6 of core.cc's 7 functions. core.cc is compiled for
+  # real here (not faked) because cft_dev_runtime.cc's ncclGet*DeviceLeInfo functions call
+  # ncclTeamCft/ncclTeamCftMultimem, and faking those would just re-derive core.cc's own
+  # already-tested logic. Links the lightweight fakes/dev_runtime_fakes.cc (hookable
+  # ncclDevrInitOnce) for core.cc's sake, and the new fakes/cft_dev_runtime_fakes.cc for
+  # cft_dev_runtime.cc's own dev_runtime.cc callees (computeLsaSize, symTeamObtain,
+  # findCommAndHostWindowFromDeviceWindow, ncclDevrWinRegEnabled) -- not the real
+  # dev_runtime.cc, which neither production file needs beyond those four symbols.
+  set(TEST_MICRO_NCCL_DEVICE_CFT_SOURCE_FILES
+    nccl-cft-dev-runtime-test.cc
+    fakes/cft_dev_runtime_fakes.cc
+    fakes/bootstrap_stubs.cc
+    fakes/topo_stubs.cc
+    fakes/transport_stubs.cc
+    fakes/nccl_stubs.cc
+    fakes/nccl_fakes.cc
+    fakes/hip_fakes.cc
+    fakes/collectives_fakes.cc
+    fakes/comm_fakes.cc
+    fakes/dev_runtime_fakes.cc
+    fakes/env_fakes.cc
+    fakes/os_fakes.cc
+    fakes/recorder_fakes.cc
+    fakes/strongstream_stubs.cc
+    fakes/tuning_fakes.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/argcheck.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/archinfo.cc
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/utils.cc
+    ../common/main_altrsmi.cpp
+    PARENT_SCOPE)
+
   # rccl-UnitTestsMicroSymKernels -- REAL src/sym_kernels.cc + generated host TU.
   set(TEST_MICRO_SYM_KERNELS_SOURCE_FILES
     sym-kernels-test.cc
@@ -1055,6 +1088,17 @@ rccl_add_micro_binary(rccl-UnitTestsMicroClientSupport
     CLIENT_SUPPORT_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/ras/client_support.cc"
   CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml")
 
+# ---------------------------------------------------------------------------
+# rccl-UnitTestsMicroNcclDeviceCft
+# ---------------------------------------------------------------------------
+rccl_add_micro_binary(rccl-UnitTestsMicroNcclDeviceCft
+  SOURCES ${TEST_MICRO_NCCL_DEVICE_CFT_SOURCE_FILES}
+  EXTRA_DEFS
+    ${_MICRO_MODE_DEFS}
+    CORE_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/nccl_device/core.cc"
+    CFT_DEV_RUNTIME_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/cft_dev_runtime.cc"
+  CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml")
+
 # ---------------------------------------------------------------------------
 # rccl-UnitTestsMicroSymKernels
 # ---------------------------------------------------------------------------

@@ -248,6 +248,32 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary:
   guard-less positional pair `libc_seam.h` / `libc_seam_undef.h` must bracket
   the production source include, after system headers and before test
   helpers. CMake pins the copied `ncclFuncStr` table to `src/init.cc`.
+- **`rccl-UnitTestsMicroNcclDeviceCft`**: `src/nccl_device/core.cc` (via
+  `CORE_CC_PATH`) + `src/cft_dev_runtime.cc` (via `CFT_DEV_RUNTIME_CC_PATH`,
+  from `nccl-cft-dev-runtime-test.cc`); suites `NcclCftDevRuntimeMicrotest.*`,
+  `NcclGetMultimemDeviceLeInfoTest.*`, `NcclGetCftDeviceLeInfoTest.*`,
+  `NcclGetPeerDeviceLeInfoTest.*` (all caught by the `Nccl*` pattern). Its own
+  binary because `rccl-UnitTestsMicro`'s `fakes/dev_runtime_micro_fakes.cc`,
+  `fakes/devcomm_fakes.cc`, and `gin-host-test.cc` between them fake 6 of
+  core.cc's 7 functions. core.cc is compiled for real here too (not faked)
+  because
+  `ncclGetMultimemDeviceLeInfo`/`ncclGetCftDeviceLeInfo`/`ncclGetPeerDeviceLeInfo`
+  call `ncclTeamCft`/`ncclTeamCftMultimem`. Links `fakes/dev_runtime_fakes.cc`
+  for core.cc's `ncclDevrInitOnce`, and the new `fakes/cft_dev_runtime_fakes.cc`
+  for `cft_dev_runtime.cc`'s own `dev_runtime.cc` callees (`computeLsaSize`,
+  `symTeamObtain`, `findCommAndHostWindowFromDeviceWindow`,
+  `ncclDevrWinRegEnabled`) -- not the real `dev_runtime.cc`, which neither
+  production file needs beyond those four symbols. This build's `CUDA_VERSION`
+  is below 13030, so the CUDA logical-endpoint driver calls throughout
+  `cft_dev_runtime.cc` are compiled out (no ROCm equivalent exists); the tests
+  cover what remains: the zero-and-succeed body of `ncclGpuCftSupport`, the
+  host-side branch logic of `computeCftSize`/`computeCftMcSize`, the
+  condition-evaluation (now-empty) bodies of `symBindTeamLe`/`symUnbindTeamLe`,
+  the trivial-failure stand-ins for `symTeamObtainUcLe`/`symTeamObtainMcLe`,
+  and the three unconditional `ncclGet*DeviceLeInfo` query functions. Lines are
+  100% covered; regions/branches have two residual, well-understood gaps
+  documented in AICOMRCCL-2825 PR4's description (`ncclParamCftEnable`, and
+  `NCCLCHECK`'s own macro-internal branches).
 - **`rccl-UnitTestsMicroSymKernels`** — the REAL `src/sym_kernels.cc` (via
   `SYM_KERNELS_CC_PATH`, from `sym-kernels-test.cc`), compiled together with the
   GENERATED `sym_kernels_host.cc` it calls into; suites `SymKernelMicrotest.*`,

@@ -0,0 +1,68 @@
+/*************************************************************************
+ * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+// See cft_dev_runtime_fakes.h.
+
+#include "cft_dev_runtime_fakes.h"
+
+#include "dev_runtime_internal.h"
+#include "fail_loud.h"
+#include "signature-drift.h"
+
+ASSERT_HOOK_MATCHES_PROD(g_computeLsaSize, computeLsaSize);
+ASSERT_HOOK_MATCHES_PROD(g_devrWinRegEnabled, ncclDevrWinRegEnabled);
+ASSERT_HOOK_MATCHES_PROD(g_findCommAndHostWindowFromDeviceWindow, findCommAndHostWindowFromDeviceWindow);
+ASSERT_HOOK_MATCHES_PROD(g_symTeamObtain, symTeamObtain);
+#undef ASSERT_HOOK_MATCHES_PROD
+
+static int DefaultComputeLsaSize(struct ncclComm*) { return 1; }
+std::function<int(struct ncclComm*)> g_computeLsaSize = DefaultComputeLsaSize;
+int computeLsaSize(struct ncclComm* comm) { return g_computeLsaSize(comm); }
+
+static bool DefaultDevrWinRegEnabled(int, ncclDevrRegisterCapability) { return true; }
+std::function<bool(int, ncclDevrRegisterCapability)> g_devrWinRegEnabled = DefaultDevrWinRegEnabled;
+bool ncclDevrWinRegEnabled(int winFlags, ncclDevrRegisterCapability capability) {
+  return g_devrWinRegEnabled(winFlags, capability);
+}
+
+ncclComm_t g_foundComm = nullptr;
+struct ncclDevrWindow* g_foundWinHost = nullptr;
+static ncclResult_t FailLoudFindCommAndHostWindow(ncclWindow_t, ncclComm_t*, struct ncclDevrWindow**) {
+  FailLoudUnfaked("cft_dev_runtime_fakes", "findCommAndHostWindowFromDeviceWindow");
+}
+std::function<ncclResult_t(ncclWindow_t, ncclComm_t*, struct ncclDevrWindow**)>
+    g_findCommAndHostWindowFromDeviceWindow = FailLoudFindCommAndHostWindow;
+ncclResult_t findCommAndHostWindowFromDeviceWindow(ncclWindow_t devWindow, ncclComm_t* foundComm,
+                                                   struct ncclDevrWindow** hostWindow) {
+  return g_findCommAndHostWindowFromDeviceWindow(devWindow, foundComm, hostWindow);
+}
+static ncclResult_t ResolveFoundCommAndWindowImpl(ncclWindow_t, ncclComm_t* foundComm,
+                                                  struct ncclDevrWindow** hostWindow) {
+  *foundComm = g_foundComm;
+  *hostWindow = g_foundWinHost;
+  return ncclSuccess;
+}
+void UseFoundCommAndWindow() { g_findCommAndHostWindowFromDeviceWindow = ResolveFoundCommAndWindowImpl; }
+
+static ncclResult_t FailLoudSymTeamObtain(struct ncclComm*, struct ncclTeam, bool, bool, bool, bool,
+                                          struct ncclDevrTeam**, bool*) {
+  FailLoudUnfaked("cft_dev_runtime_fakes", "symTeamObtain");
+}
+std::function<ncclResult_t(struct ncclComm*, struct ncclTeam, bool, bool, bool, bool, struct ncclDevrTeam**, bool*)>
+    g_symTeamObtain = FailLoudSymTeamObtain;
+ncclResult_t symTeamObtain(struct ncclComm* comm, struct ncclTeam team, bool multimem, bool counted, bool wantsLeUc,
+                           bool wantsLeMc, struct ncclDevrTeam** outTeam, bool* needBarrier) {
+  return g_symTeamObtain(comm, team, multimem, counted, wantsLeUc, wantsLeMc, outTeam, needBarrier);
+}
+
+void ResetCftDevRuntimeFakes() {
+  g_computeLsaSize = DefaultComputeLsaSize;
+  g_devrWinRegEnabled = DefaultDevrWinRegEnabled;
+  g_findCommAndHostWindowFromDeviceWindow = FailLoudFindCommAndHostWindow;
+  g_foundComm = nullptr;
+  g_foundWinHost = nullptr;
+  g_symTeamObtain = FailLoudSymTeamObtain;
+}
@@ -0,0 +1,51 @@
+/*************************************************************************
+ * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+// Seams for src/cft_dev_runtime.cc's callees that live in src/dev_runtime.cc
+// (computeLsaSize, symTeamObtain, findCommAndHostWindowFromDeviceWindow) and
+// src/dev_runtime.cc's ncclDevrWinRegEnabled. Targets linking this file
+// compile the real cft_dev_runtime.cc but not the much larger dev_runtime.cc,
+// so these four symbols need a seam of their own rather than the full file.
+
+#ifndef RCCL_TEST_HOST_CFT_DEV_RUNTIME_FAKES_H_
+#define RCCL_TEST_HOST_CFT_DEV_RUNTIME_FAKES_H_
+
+#include <functional>
+
+#include "nccl.h"
+// ncclDevrRegisterCapability is a plain (non-fixed-underlying-type) enum, which C++
+// does not allow forward-declaring -- pull in the real declarations instead.
+#include "dev_runtime.h"
+
+struct ncclComm;
+struct ncclTeam;
+
+// computeLsaSize: generous default (1) -- matches a comm with no LSA peers.
+extern std::function<int(struct ncclComm*)> g_computeLsaSize;
+
+// ncclDevrWinRegEnabled: default true -- matches a window registered for every capability.
+extern std::function<bool(int, ncclDevrRegisterCapability)> g_devrWinRegEnabled;
+
+// findCommAndHostWindowFromDeviceWindow: fail-loud default; a test must set
+// g_foundComm/g_foundWinHost and call UseFoundCommAndWindow() (or install its own
+// hook) before calling a function that depends on it.
+extern std::function<ncclResult_t(ncclWindow_t, ncclComm_t*, struct ncclDevrWindow**)>
+    g_findCommAndHostWindowFromDeviceWindow;
+extern ncclComm_t g_foundComm;
+extern struct ncclDevrWindow* g_foundWinHost;
+// Installs the default hook that returns g_foundComm/g_foundWinHost directly.
+void UseFoundCommAndWindow();
+
+// symTeamObtain: fail-loud default; a test must install its own hook (e.g. one
+// that heap-allocates a struct ncclDevrTeam with a populated worldRankList)
+// before calling ncclGetMultimemDeviceLeInfo.
+extern std::function<ncclResult_t(struct ncclComm*, struct ncclTeam, bool, bool, bool, bool, struct ncclDevrTeam**,
+                                   bool*)>
+    g_symTeamObtain;
+
+void ResetCftDevRuntimeFakes();
+
+#endif // RCCL_TEST_HOST_CFT_DEV_RUNTIME_FAKES_H_