test(rccl): add host-only unit tests for src/transport/net.cc

projects/rccl/test/host/CMakeLists.txt

@@ -13,6 +13,12 @@ rccl_assert_source_line("${_rccl_src_root}/graph/connect.cc"
   "NCCL_PARAM(MaxNchannels, \"MAX_NCHANNELS\", -2)" "fakes/tuning_fakes.cc: ncclParamMaxNchannels()")
 rccl_assert_source_line("${_rccl_src_root}/enqueue/enqueue.cc"
   "RCCL_PARAM(ForceCe, \"FORCE_CE\", 1)" "fakes/wrap_fakes.cc: rcclParamForceCe()")
+rccl_assert_source_line("${_rccl_src_root}/misc/rocmwrap.cc"
+  "RCCL_PARAM(ForceEnableDMABUF, \"FORCE_ENABLE_DMABUF\", 0);"
+  "fakes/rocmwrap_fakes.cc: rcclParamForceEnableDMABUF()")
+rccl_assert_source_line("${CMAKE_CURRENT_LIST_DIR}/fakes/rocmwrap_fakes.cc"
+  "int64_t rcclParamForceEnableDMABUF() { return g_loadParam(\"RCCL_FORCE_ENABLE_DMABUF\", 0); }"
+  "the ForceEnableDMABUF fake default")
 # src/enqueue/task_prep/*.cc declare these five but define none, and param_redirect.h only rewrites
 # NCCL_PARAM where it expands, so the hand-written bodies below are the only definitions on that link line.
 rccl_assert_source_line("${_rccl_src_root}/init.cc"
@@ -880,6 +886,30 @@ function(rccl_define_micro_source_lists)
     ../common/main_altrsmi.cpp
     PARENT_SCOPE)
 
+  # rccl-UnitTestsMicroNet -- src/transport/net.cc. Own binary: transport_stubs.cc, which most
+  # targets link, defines rcclUseAinic and netTransport, and net.cc defines both.
+  set(TEST_MICRO_NET_SOURCE_FILES
+    net-test.cc
+    fakes/collectives_fakes.cc           # func/algo/proto name tables
+    fakes/env_fakes.cc
+    fakes/hip_fakes.cc
+    fakes/nccl_fakes.cc
+    fakes/nccl_stubs.cc                  # core floor, incl. proxy profiler hooks
+    fakes/net_ib_cast_common_fakes.cc    # netIbCast
+    fakes/net_ib_cast_connect_fakes.cc   # rcclCastNetP2pPolicy
+    fakes/os_fakes.cc
+    fakes/proxy_fakes.cc
+    fakes/rccl_wrap_fakes.cc             # NET registration segment count / commit
+    fakes/rocmwrap_fakes.cc
+    fakes/topo_stubs.cc
+    fakes/transport_net_ib_connect_net_fakes.cc  # rcclNetP2pPolicy, GDR flush-flag accessors
+    fakes/transport_net_ib_fakes.cc
+    fakes/transport_p2p_fakes.cc         # src/transport/p2p.cc shareable buffers
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/archinfo.cc  # real IsArchMatch
+    ${RCCL_HOST_HIPIFY_DIR}/src/misc/utils.cc     # real parseStringList/matchIfList for NCCL_IB_HCA
+    ../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
@@ -2293,6 +2323,22 @@ rccl_add_micro_binary(rccl-UnitTestsMicroNetIbCastInit
     NET_IB_CAST_GID_INSPECT_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/transport/net_ib_cast/net_ib_cast_gid_inspect.cc"
   CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_net_ib_cast_init.yaml")
 
+# ---------------------------------------------------------------------------
+# rccl-UnitTestsMicroNet
+# ---------------------------------------------------------------------------
+rccl_add_micro_binary(rccl-UnitTestsMicroNet
+  SOURCES ${TEST_MICRO_NET_SOURCE_FILES}
+  EXTRA_INCLUDES
+    ${RCCL_HOST_HIPIFY_DIR}/src/transport/net_ib_cast  # the net_ib_cast fakes' headers
+    ${RCCL_HOST_HIPIFY_DIR}/src/transport              # net_ib_limits.h
+  EXTRA_DEFS
+    ${_MICRO_MODE_DEFS}
+    ${_MICRO_DMABUF_EXPORT_DEF}  # standalone: the cuMem DMA-BUF arms in-build gets from rccl
+    # net_tmp.cc, not net.cc: src/CMakeLists.txt lists plugin/net.cc first, so
+    # add_file_unique() renames this one. Re-check after any source-list reorder.
+    NET_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/transport/net_tmp.cc"
+  CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_net.yaml")
+
 # =====================================================================
 # Standalone-only tail: coverage + no CTest registration.
 # =====================================================================

projects/rccl/test/host/MICROTEST_README.md

@@ -520,6 +520,10 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary:
   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`.
+- **`rccl-UnitTestsMicroNet`**: `src/transport/net.cc` (via `NET_CC_PATH`, suites
+  `Net*Microtest.*`), driven through a fake `ncclNet_t`. Its own binary: `transport_stubs.cc`,
+  which most targets link, defines `rcclUseAinic` and `netTransport`, and `net.cc` defines
+  both. See `test_categories_micro_net.yaml`.
 
 - **`rccl-UnitTestsMicroGinHostFinalize`**: `src/gin/gin_host.cc` (via
   `GIN_HOST_CC_PATH`, suite `GinHostFinalizeTest.*`): host finalize, connect
@@ -732,13 +736,14 @@ symbol.
 | `src/rma/*.cc` | `fakes/rma_fakes.cc` |
 | `src/scheduler/*.cc`'s own public entry points (targets that don't compile the real files, e.g. `rccl-UnitTestsMicroEnqueue`) and the deep launch paths | `fakes/sched_stubs.cc` |
 | `src/sym_kernels.cc` | `fakes/sym_kernels_fakes.cc` |
-| `src/transport/net_ib/connect.cc` params (`ncclParamIbQpsPerConn`) | `fakes/transport_net_ib_fakes.cc` |
+| `src/transport/net_ib/` (`ncclParamIbQpsPerConn`, `ncclIbGidInfoQuery`, `ncclNetIb`, GDR flush-flag support and enable) | `fakes/transport_net_ib_fakes.cc` |
+| `src/transport/net_ib/connect.cc`'s `rcclNetP2pPolicy` and GDR flush-flag accessors (targets that do not compile the real file) | `fakes/transport_net_ib_connect_net_fakes.cc` |
 | `src/transport/net_ib_cast/common.cc` (targets that do not compile the real file) | `fakes/net_ib_cast_common_fakes.cc` |
 | `src/transport/net_ib_cast/{connect,p2p,p2p_resiliency,p2p_resiliency_recovery,qp_sharing,scheduler}.cc` (targets that do not compile the real file) | `fakes/net_ib_cast_<file>_fakes.cc` |
 | `src/transport/net_ib_cast/init.cc` (targets that do not compile the real file) | `fakes/net_ib_cast_init_fakes.cc` |
 | generated `sym_kernels_host.cc` (targets that do not compile it) | `fakes/sym_kernels_index_fakes.cc` |
 | rocSHMEM host API | `fakes/rocshmem_fakes.cc` |
-| `src/transport/p2p.cc` shareable-buffer entry points (`rccl-UnitTestsMicroDiagnostics`) | `fakes/transport_p2p_fakes.cc` |
+| `src/transport/p2p.cc` shareable-buffer entry points (`rccl-UnitTestsMicroDiagnostics`, `rccl-UnitTestsMicroNet`) | `fakes/transport_p2p_fakes.cc` |
 | `src/transport/*`, `src/plugin/net.cc` | `fakes/transport_stubs.cc` |
 | libc (`gethostname`, `dladdr`), process-wide for units linking libc directly | `fakes/libc_interposers.cc` |
 | libc macro-renamed in one unit via `fakes/libc_seam.h` (sockets/stdio/exit for `ras/client.cc` and `ras/client_support.cc`; `gethostname`/`access` for `diagnostics/ib_write_bw.cc`; heap for `graph/rccl_graph_gen.cc`) | `fakes/libc_fakes.cc` |
@@ -1079,7 +1084,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-UnitTestsMicroGinHostFinalize`, `rccl-UnitTestsMicroDdaIpc`, `rccl-UnitTestsMicroCe`, `rccl-UnitTestsMicroDiagnostics`,
-`rccl-UnitTestsMicroGraph`, `rccl-UnitTestsMicroDiagnostics`, `rccl-UnitTestsMicroDda`,
+`rccl-UnitTestsMicroGraph`, `rccl-UnitTestsMicroNet`, `rccl-UnitTestsMicroDiagnostics`, `rccl-UnitTestsMicroDda`,
 `rccl-UnitTestsMicroSymKernels`, `rccl-UnitTestsMicroTaskPrep`, `rccl-UnitTestsMicroNvls`,
 `rccl-UnitTestsMicroNetIbCastDataPath`, `rccl-UnitTestsMicroRmaProxy` and `rccl-UnitTestsMicroNetIbCastInit` — **without configuring/building all of
 librccl**. It compiles just the tests + fakes + the hipified unit-under-test
@@ -1155,6 +1160,7 @@ cmake --build build -j"$(nproc)"
 ./build/rccl-UnitTestsMicroProxy              # src/proxy.cc tests
 ./build/rccl-UnitTestsMicroDda                # src/algorithms/dda/ tests
 ./build/rccl-UnitTestsMicroNetIbCastDataPath  # net_ib_cast p2p/scheduler/ops-fault tests
+./build/rccl-UnitTestsMicroNet                # src/transport/net.cc tests
 ./build/rccl-HostUnitTests
 ./build/rccl-UnitTestsMicroRoctxNvml          # misc/roctx.cc + misc/nvmlwrap_stub.cc tests
 ./build/rccl-UnitTestsMicroRoctxNvml-roctx    # same, ROCTX_ENABLE arm

projects/rccl/test/host/fakes/hip_fakes.cc

@@ -602,6 +602,7 @@ static hipError_t DefaultHipStreamBatchMemOp(hipStream_t, unsigned int,
 std::function<hipError_t(hipStream_t, unsigned int, hipStreamBatchMemOpParams*, unsigned int)>
     g_hipStreamBatchMemOp = DefaultHipStreamBatchMemOp;
 
+
 #ifdef CE_BATCH_ASYNC_SUPPORTED
 static hipError_t DefaultHipMemcpyBatchAsync(void**, void**, size_t*, size_t, hipMemcpyAttributes*, size_t*,
                                              size_t, size_t*, hipStream_t)
@@ -1016,6 +1017,7 @@ hipError_t hipStreamWriteValue32(hipStream_t, void*, std::uint32_t, unsigned int
     FailLoudUnfaked("hip_fakes", "hipStreamWriteValue32");
 }
 
+
 // ce_coll.cc's batch path; in-RCCL builds inherit CE_BATCH_ASYNC_SUPPORTED, standalone ones compile the call out.
 #ifdef CE_BATCH_ASYNC_SUPPORTED
 hipError_t hipMemcpyBatchAsync(void** dsts, void** srcs, size_t* sizes, size_t count, hipMemcpyAttributes* attrs,

projects/rccl/test/host/fakes/nccl_fakes.cc

@@ -57,6 +57,8 @@ ASSERT_HOOK_MATCHES_PROD(g_memTrackImportFromPeer,    ncclMemTrackImportFromPeer
 ASSERT_HOOK_MATCHES_PROD(g_memTrack,                  ncclMemTrack);
 ASSERT_HOOK_MATCHES_PROD(g_memUntrackDynamic,         ncclMemUntrackDynamic);
 ASSERT_HOOK_MATCHES_PROD(g_dynMemMarkExportToPeer,    ncclDynMemMarkExportToPeer);
+ASSERT_HOOK_MATCHES_PROD(g_shmImportShareableBuffer,  ncclShmImportShareableBuffer);
+ASSERT_HOOK_MATCHES_PROD(g_shmIpcClose,               ncclShmIpcClose);
 
 #undef ASSERT_HOOK_MATCHES_PROD
 
@@ -241,20 +243,32 @@ ncclResult_t ncclRegLocalIsValid(struct ncclReg* reg, bool* isValid)
     return g_regLocalIsValid(reg, isValid);
 }
 
-ncclResult_t ncclShmImportShareableBuffer(struct ncclComm*  /*comm*/,
-                                          int               /*proxyRank*/,
-                                          ncclShmIpcDesc_t* /*desc*/,
-                                          void**            /*hptr*/,
-                                          void**            /*dptr*/,
-                                          ncclShmIpcDesc_t* /*descOut*/)
+static ncclResult_t DefaultShmImportShareableBuffer(struct ncclComm*, int, ncclShmIpcDesc_t*, void**, void**,
+                                                    ncclShmIpcDesc_t*)
 {
     return ncclSystemError;
 }
+std::function<ncclResult_t(struct ncclComm*, int, ncclShmIpcDesc_t*, void**, void**, ncclShmIpcDesc_t*)>
+    g_shmImportShareableBuffer = DefaultShmImportShareableBuffer;
+ncclResult_t ncclShmImportShareableBuffer(struct ncclComm*  comm,
+                                          int               proxyRank,
+                                          ncclShmIpcDesc_t* desc,
+                                          void**            hptr,
+                                          void**            dptr,
+                                          ncclShmIpcDesc_t* descOut)
+{
+    return g_shmImportShareableBuffer(comm, proxyRank, desc, hptr, dptr, descOut);
+}
 
-ncclResult_t ncclShmIpcClose(ncclShmIpcDesc_t* /*desc*/)
+static ncclResult_t DefaultShmIpcClose(ncclShmIpcDesc_t*)
 {
     return ncclSuccess;
 }
+std::function<ncclResult_t(ncclShmIpcDesc_t*)> g_shmIpcClose = DefaultShmIpcClose;
+ncclResult_t ncclShmIpcClose(ncclShmIpcDesc_t* desc)
+{
+    return g_shmIpcClose(desc);
+}
 
 // src/plugin/profiler.cc, which has no fakes file yet. plugin/net.cc only hands its address to a plugin's init.
 ncclResult_t ncclProfilerCallback(void**, int, void*, int64_t, void*)
@@ -629,6 +643,8 @@ void ResetNcclFakes()
     g_ncclTopoCheckNet             = DefaultTopoCheckNet;
     g_regLocalIsValid              = DefaultRegLocalIsValid;
     g_shmAllocateShareableBuffer   = DefaultShmAllocateShareableBuffer;
+    g_shmImportShareableBuffer     = DefaultShmImportShareableBuffer;
+    g_shmIpcClose                  = DefaultShmIpcClose;
     g_commGraphRegister            = DefaultCommGraphRegister;
     g_commGraphDeregister          = DefaultCommGraphDeregister;
     g_ncclSetThreadName            = DefaultSetThreadName;

projects/rccl/test/host/fakes/nccl_fakes.h

@@ -20,6 +20,10 @@
 #include <cstddef>
 #include <cstdint>
 #include <functional>
+#include <initializer_list>
+#include <string>
+#include <utility>
+#include <vector>
 
 #include "nccl.h"
 #include "strongstream.h"
@@ -28,6 +32,7 @@
 #include <hip/hip_runtime.h>
 
 struct ncclReg;
+struct shmIpcDesc;  // ncclShmIpcDesc_t
 
 // ncclStrongStreamAcquire: by default returns ncclSuccess with *stream=nullptr
 // (matching the stub's prior behaviour). Tests that need to exercise the
@@ -124,6 +129,18 @@ extern std::function<ncclResult_t(struct ncclMemManager*, void*, struct ncclMemU
 extern std::function<int64_t(const char* /*env*/, int64_t /*deftVal*/)>
     g_loadParam;
 
+// A g_loadParam that answers the listed env names and leaves the rest at their defaults.
+inline std::function<int64_t(const char*, int64_t)> ParamOverride(
+    std::initializer_list<std::pair<const char*, int64_t>> kv) {
+  std::vector<std::pair<std::string, int64_t>> table(kv.begin(), kv.end());
+  return [table](const char* env, int64_t deft) {
+    for (const auto& [name, value] : table) {
+      if (name == env) return value;
+    }
+    return deft;
+  };
+}
+
 // 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 /*dev1*/, int /*dev2*/, bool* /*isXGMI*/, int /*maxInter*/)>
@@ -168,6 +185,15 @@ extern std::function<ncclResult_t(size_t /*size*/, bool /*legacy*/,
 // ncclSetThreadName(pthread_t, fmt, ...) hands the formatted name here.
 extern std::function<void(pthread_t, const char* /*name*/)> g_ncclSetThreadName;
 
+// ncclShmImportShareableBuffer / ncclShmIpcClose: a transport maps (and later
+// closes) a host buffer a proxy in another process created. Import fails by
+// default; close succeeds.
+extern std::function<ncclResult_t(struct ncclComm*, int /*proxyRank*/,
+                                  struct shmIpcDesc* /*desc*/, void** /*hptr*/,
+                                  void** /*dptr*/, struct shmIpcDesc* /*descOut*/)>
+    g_shmImportShareableBuffer;
+extern std::function<ncclResult_t(struct shmIpcDesc*)> g_shmIpcClose;
+
 // Restore every NCCL controllable seam in this header to its default.
 // Called by ResetP2pFakes(); exposed for tests that only touch NCCL hooks.
 // Hands back an fd the caller must close. Defaults to ncclSystemError (the

projects/rccl/test/host/fakes/nccl_stubs.cc

@@ -143,6 +143,19 @@ uint8_t ncclProfilerDeviceMode(int eActivationMask) {
   return mode;
 }
 bool ncclProfilerProxyDiagEnabled(void) { return false; }
+ncclResult_t ncclProfilerStartSendProxyStepEvent(int, struct ncclProxyArgs*, int) { return ncclSuccess; }
+ncclResult_t ncclProfilerStartRecvProxyStepEvent(int, struct ncclProxyArgs*, int) { return ncclSuccess; }
+ncclResult_t ncclProfilerStopProxyStepEvent(int, struct ncclProxyArgs*, int) { return ncclSuccess; }
+ncclResult_t ncclProfilerStopProxyOpEvent(int, struct ncclProxyArgs*) { return ncclSuccess; }
+ncclResult_t ncclProfilerRecordProxyOpEventState(int, struct ncclProxyArgs*, ncclProfilerEventState_t) {
+  return ncclSuccess;
+}
+ncclResult_t ncclProfilerRecordProxyStepEventState(int, struct ncclProxyArgs*, int, ncclProfilerEventState_t) {
+  return ncclSuccess;
+}
+ncclResult_t ncclProfilerRecordProxyDiagState(int, struct ncclProxyArgs*, uint8_t, int64_t, int64_t, uint16_t) {
+  return ncclSuccess;
+}
 ncclResult_t ncclRasCommFini(const struct ncclComm* comm) { return ncclSuccess; }
 ncclResult_t ncclRunRasDiagnostics(struct ncclComm* comm) { return ncclSuccess; }
 ncclResult_t ncclRunDiagnostics(struct ncclComm* comm) { return ncclSuccess; }

projects/rccl/test/host/fakes/net_ib_cast_common_fakes.cc

@@ -22,6 +22,8 @@ struct ncclIbDev IbCastDevs[MAX_IB_DEVS];
 int IbCastRelaxedOrderingEnabled = 0;
 ncclProfilerCallback_t IbCastProfilerFunction;
 std::thread IbCastAsyncThread;
+// net.cc compares a plugin's address against it.
+ncclNet_t netIbCast = {"IB-CAST"};
 
 int64_t ncclParamIbCastSplitDataOnQps() { return g_loadParam("IB_SPLIT_DATA_ON_QPS", 0); }
 int64_t ncclParamIbCastPrepostReceiveWorkRequests() { return g_loadParam("IB_PREPOST_RECEIVE_WORK_REQUESTS", -2); }

projects/rccl/test/host/fakes/net_ib_cast_connect_fakes.cc

@@ -119,6 +119,7 @@ ASSERT_HOOK_MATCHES_PROD(g_ibCastCloseListen, IbCastCloseListen);
 ASSERT_HOOK_MATCHES_PROD(g_ibCastGetGidIndex, IbCastGetGidIndex);
 ASSERT_HOOK_MATCHES_PROD(g_ibCastGidInfoQuery, IbCastGidInfoQuery);
 ASSERT_HOOK_MATCHES_PROD(g_ibCastSetTrafficClass, IbCastSetTrafficClass);
+ASSERT_HOOK_MATCHES_PROD(g_rcclCastNetP2pPolicy, rcclCastNetP2pPolicy);
 
 ncclResult_t IbCastQpCreate(struct ncclIbQp* qp, struct ncclIbQpCreateAttr* attr) { return g_ibCastQpCreate(qp, attr); }
 void IbCastBuildDataQpCreateAttr(struct ncclIbNetCommBase* base, int devIndex, struct ncclIbQpCreateAttr* out) {
@@ -161,6 +162,12 @@ ncclResult_t IbCastGidInfoQuery(struct ibv_context* context, uint8_t portNum, st
 void IbCastSetTrafficClass(void* ctx, int trafficClass) { g_ibCastSetTrafficClass(ctx, trafficClass); }
 void IbCastReportMatchingScheme() { g_ibCastReportMatchingSchemeCalls++; }
 
+static ncclResult_t DefaultCastNetP2pPolicy(void*, int) {
+  FailLoudUnfaked("net_ib_cast_connect_fakes", "rcclCastNetP2pPolicy");
+}
+std::function<ncclResult_t(void*, int)> g_rcclCastNetP2pPolicy = DefaultCastNetP2pPolicy;
+ncclResult_t rcclCastNetP2pPolicy(void* handle, int isP2p) { return g_rcclCastNetP2pPolicy(handle, isP2p); }
+
 void ResetNetIbCastConnectFakes() {
   g_ibCastQpCreate = DefaultIbCastQpCreate;
   g_ibCastBuildDataQpCreateAttr = DefaultIbCastBuildDataQpCreateAttr;
@@ -182,4 +189,5 @@ void ResetNetIbCastConnectFakes() {
   g_ibCastGidInfoQuery = DefaultIbCastGidInfoQuery;
   g_ibCastSetTrafficClass = DefaultIbCastSetTrafficClass;
   g_ibCastReportMatchingSchemeCalls = 0;
+  g_rcclCastNetP2pPolicy = DefaultCastNetP2pPolicy;
 }

projects/rccl/test/host/fakes/net_ib_cast_connect_fakes.h

@@ -47,6 +47,8 @@ extern std::function<ncclResult_t(struct ibv_context*, uint8_t, struct ibv_port_
     g_ibCastGidInfoQuery;
 extern std::function<void(void*, int)> g_ibCastSetTrafficClass;
 extern int g_ibCastReportMatchingSchemeCalls;
+// The P2P QP policy net.cc applies to an IB-cast connect handle. Defaults to fail-loud.
+extern std::function<ncclResult_t(void* /*handle*/, int /*isP2p*/)> g_rcclCastNetP2pPolicy;
 
 void ResetNetIbCastConnectFakes();