net-ib: refactor fault injection tests

Fix fault injection API issues from review:
- ncclIbCastFaultSetQpDelay/SetQpError: use comm->base.nqps for bounds
  check instead of NCCL_IB_MAX_QPS
- ncclIbCastFaultClear: atomically reset fatalErrorCount to 0
- move net_ib_fault_inject.h into transport/net_ib_cast/; drop local
  #define NCCL_IB_MAX_QPS 128, include net_ib_cast_inspect.h and add
  static_assert so a size mismatch becomes a compile error
- fault hook in IbCastMultiSend: use IbCastStatsFatalError (renamed from
  ncclIbStatsFatalError in asanniko's split)
- FaultInjCastQpErrorClearRecovers: ASSERT_EQ on SetQpError return value;
  drain recvReq before DeregisterMemory
- FaultInjCastSingleQpErrorIsFatal: EXPECT_EQ on ncclIbCastSetTokens and
  ncclIbCastFaultClear return values

Consolidate NCCL_IB_MAX_QPS definition and reorganize headers:
- add src/transport/net_ib_limits.h as single source of truth for
  NCCL_IB_MAX_QPS 128, removing duplicate defines from net_ib.cc,
  common_cast.h, and net_ib_cast_inspect.h
- move net_ib_cast_inspect.h from src/include/ to src/transport/net_ib_cast/
  alongside net_ib_fault_inject.h; update CMakeLists include paths
- group net_ib_limits.h, net_ib_cast_inspect.h, net_ib_fault_inject.h
  together next to transport/net_ib.cc in src/CMakeLists.txt

Reduce test code duplication:
- add helpers to NetIbMPITestBase: DrainRecvRequest, TeardownConnection,
  ExpectEqualWeightInitTokens, ExpectActiveTokenSumInvariant
- use them in FaultInjectTests.cpp and CastTests.cpp to eliminate ~300
  lines of repeated teardown and sched-state assertion code

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>

Author: i-kosarev

projects/rccl/src/CMakeLists.txt

@@ -115,7 +115,6 @@ set(SRC_FILES
   include/nccl_device.h
   include/net_device.h
   include/net.h
-  include/net_ib_cast_inspect.h
   include/nvmlwrap.h
   include/nvtx.h
   include/nvtx_payload_schemas.h
@@ -219,7 +218,6 @@ set(SRC_FILES
   include/plugin/net/net_v11.h
   include/plugin/profiler/net_ib_v1.h
   include/plugin/profiler/net_ib.h
-  include/net_ib_fault_inject.h
   include/plugin/profiler/net_socket_v1.h
   include/plugin/profiler/net_socket.h
   include/plugin/profiler/profiler_v1.h
@@ -296,6 +294,9 @@ set(SRC_FILES
   transport/generic.cc
   transport/net.cc
   transport/net_ib.cc
+  transport/net_ib_limits.h
+  transport/net_ib_cast/net_ib_cast_inspect.h
+  transport/net_ib_cast/net_ib_fault_inject.h
   # net_ib_rocm.cc is generated by rocmIb.cmake directly into the hipify staging area
   # so it is not listed here - it's added to HIP_SOURCES separately below
   transport/net_socket.cc

projects/rccl/src/transport/net_ib.cc

@@ -1164,7 +1164,7 @@ ncclResult_t ncclIbGetProperties(int dev, ncclNetProperties_t* props) {
 #define MAX_REQUESTS (NCCL_NET_MAX_REQUESTS*NCCL_NET_IB_MAX_RECVS)
 static_assert(MAX_REQUESTS <= 256, "request id are encoded in wr_id and we need up to 8 requests ids per completion");
 
-#define NCCL_IB_MAX_QPS 128
+#include "net_ib_limits.h"
 
 // Per-QP connection metatdata
 struct ncclIbQpInfo {

projects/rccl/src/transport/net_ib_cast/CMakeLists.txt

@@ -1,5 +1,7 @@
 # Transport sources (paths relative to PROJECT_SOURCE_DIR to match SRC_FILES convention)
 set(NET_IB_CAST_SOURCES
+    src/transport/net_ib_limits.h
+    src/transport/net_ib_cast/net_ib_cast_inspect.h
     src/transport/net_ib_cast/common_cast.h
     src/transport/net_ib_cast/connect_cast.h
     src/transport/net_ib_cast/gin_cast.h
@@ -26,6 +28,7 @@ set(NET_IB_CAST_SOURCES
 # is forwarded to target_include_directories(rccl ...).
 set(NET_IB_CAST_INCLUDE_DIRS
     "${CMAKE_BINARY_DIR}/hipify/src/transport/net_ib_cast"
+    "${CMAKE_BINARY_DIR}/hipify/src/transport"
     PARENT_SCOPE)
 
 # Propagate sources to parent scope so src/CMakeLists.txt can use ${NET_IB_CAST_SOURCES}

projects/rccl/src/transport/net_ib_cast/common_cast.h

@@ -748,7 +748,6 @@ ncclResult_t IbCastSetNetAttr(void *ctx, ncclNetAttr_t *netAttr);
 #define NCCL_CTS_QP_SLOT_INVALID 0xFF
 
 // IB-CAST specific infrastructure
-#define NCCL_IB_MAX_QPS 128
 
 #define NSEC_PER_USEC           1000ULL
 #define NSEC_PER_MSEC           (NSEC_PER_USEC * 1000)

…s/rccl/src/include/net_ib_cast_inspect.h …nsport/net_ib_cast/net_ib_cast_inspect.hprojects/rccl/src/include/net_ib_cast_inspect.h renamed to projects/rccl/src/transport/net_ib_cast/net_ib_cast_inspect.h projects/rccl/src/include/net_ib_cast_inspect.h renamed to projects/rccl/src/transport/net_ib_cast/net_ib_cast_inspect.h

@@ -8,6 +8,7 @@
 #define RCCL_NET_IB_CAST_INSPECT_H_
 
 #include <stdint.h>
+#include "net_ib_limits.h"
 
 #ifdef __cplusplus
 #include "nccl.h"  /* ncclResult_t */
@@ -19,14 +20,10 @@ extern "C" {
 /*
  * Test-only introspection API for the net-ib-cast WRR scheduler.
  *
- * Implementation lives in src/transport/net_ib_cast.cc and is compiled into
- * librccl.so.  Both the library and the unit tests include this one header so
- * the struct layout and function signatures can never diverge.
- *
- * NCCL_IB_MAX_QPS is the maximum number of QPs per connection.
- * Defined here so both the library and the tests share the same value.
+ * Implementation lives in src/transport/net_ib_cast/scheduler.cc and is
+ * compiled into librccl.so.  Both the library and the unit tests include this
+ * one header so the struct layout and function signatures can never diverge.
  */
-#define NCCL_IB_MAX_QPS 128
 
 struct ncclIbCastSchedState {
   int      nqps;

…s/rccl/src/include/net_ib_fault_inject.h …nsport/net_ib_cast/net_ib_fault_inject.hprojects/rccl/src/include/net_ib_fault_inject.h renamed to projects/rccl/src/transport/net_ib_cast/net_ib_fault_inject.h projects/rccl/src/include/net_ib_fault_inject.h renamed to projects/rccl/src/transport/net_ib_cast/net_ib_fault_inject.h

@@ -22,13 +22,11 @@ extern "C" {
 /*
  * Test-only per-QP fault injection API for the net-ib CAST transport.
  *
- * Implemented in src/transport/net_ib_cast.cc (CAST multi-QP path).
- *
- * NCCL_IB_MAX_QPS is defined in net_ib_cast_inspect.h; guard against
- * double-definition when both headers are included together.
+ * Implemented in src/transport/net_ib_cast/p2p.cc (CAST multi-QP path).
  */
-#ifndef NCCL_IB_MAX_QPS
-#define NCCL_IB_MAX_QPS 128
+#include "net_ib_cast_inspect.h"
+#ifdef __cplusplus
+static_assert(NCCL_IB_MAX_QPS == 128, "fault injection arrays sized for 128 QPs; update if NCCL_IB_MAX_QPS changes");
 #endif
 
 /* ── CAST path (net_ib_cast.cc) ───────────────────────────────────────── */
@@ -39,7 +37,7 @@ extern "C" {
 ncclResult_t ncclIbCastFaultSetQpDelay(void* sendComm, int qpIdx, uint32_t delayUs);
 
 /* Arm error injection on a specific QP index.
- * When armed, the hook calls ncclIbStatsFatalError and then returns
+ * When armed, the hook calls IbCastStatsFatalError and then returns
  * ncclSystemError instead of calling wrap_ibv_post_send.
  * Set inject=false to disarm. */
 ncclResult_t ncclIbCastFaultSetQpError(void* sendComm, int qpIdx, bool inject);

projects/rccl/src/transport/net_ib_cast/p2p.cc

@@ -298,7 +298,7 @@ ncclResult_t IbCastMultiSend(struct ncclIbSendComm* comm, int slot, int nqps, in
       const uint32_t faultDelay = comm->base.faultQpDelayUs[qpIndex];
       if (faultDelay) usleep(faultDelay);
       if (comm->base.faultQpError[qpIndex]) {
-        ncclIbStatsFatalError(&comm->base.stats);
+        IbCastStatsFatalError(&comm->base.stats);
         return ncclSystemError;
       }
     }
@@ -1105,15 +1105,15 @@ ncclResult_t IbCastTest(void* request, int* done, int* sizes) {
 ncclResult_t ncclIbCastFaultSetQpDelay(void* sendComm, int qpIdx, uint32_t delayUs) {
   if (!sendComm) return ncclInvalidArgument;
   struct ncclIbSendComm* comm = (struct ncclIbSendComm*)sendComm;
-  if (qpIdx < 0 || qpIdx >= NCCL_IB_MAX_QPS) return ncclInvalidArgument;
+  if (qpIdx < 0 || qpIdx >= comm->base.nqps) return ncclInvalidArgument;
   comm->base.faultQpDelayUs[qpIdx] = delayUs;
   return ncclSuccess;
 }
 
 ncclResult_t ncclIbCastFaultSetQpError(void* sendComm, int qpIdx, bool inject) {
   if (!sendComm) return ncclInvalidArgument;
   struct ncclIbSendComm* comm = (struct ncclIbSendComm*)sendComm;
-  if (qpIdx < 0 || qpIdx >= NCCL_IB_MAX_QPS) return ncclInvalidArgument;
+  if (qpIdx < 0 || qpIdx >= comm->base.nqps) return ncclInvalidArgument;
   comm->base.faultQpError[qpIdx] = inject;
   return ncclSuccess;
 }
@@ -1123,6 +1123,7 @@ ncclResult_t ncclIbCastFaultClear(void* sendComm) {
   struct ncclIbSendComm* comm = (struct ncclIbSendComm*)sendComm;
   memset(comm->base.faultQpDelayUs, 0, sizeof(comm->base.faultQpDelayUs));
   memset(comm->base.faultQpError, 0, sizeof(comm->base.faultQpError));
+  __atomic_store_n(&comm->base.stats.fatalErrorCount, 0, __ATOMIC_RELEASE);
   return ncclSuccess;
 }
 

projects/rccl/src/transport/net_ib_cast/scheduler.cc

@@ -425,7 +425,7 @@ ncclResult_t IbCastQpSchedFreeRemap(struct ncclIbRemapWrId* r) {
 // Test introspection API — exposes internal WRR scheduler state from a
 // sendComm handle.  Only intended for unit tests; not part of the public net
 // plugin ABI.
-// Struct definition and function prototypes live in src/include/net_ib_cast_inspect.h.
+// Struct definition and function prototypes live in src/transport/net_ib_cast/net_ib_cast_inspect.h.
 // =============================================================================
 
 // ncclIbCastGetSchedState — copy WRR scheduler state out of a sendComm.

projects/rccl/src/transport/net_ib_limits.h

@@ -0,0 +1,13 @@
+/*************************************************************************
+ * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#ifndef RCCL_NET_IB_LIMITS_H_
+#define RCCL_NET_IB_LIMITS_H_
+
+/* Maximum number of QPs per IB connection (CAST and regular paths). */
+#define NCCL_IB_MAX_QPS 128
+
+#endif /* RCCL_NET_IB_LIMITS_H_ */

projects/rccl/test/CMakeLists.txt

@@ -94,6 +94,8 @@ if(BUILD_TESTS)
     ${PROJECT_BINARY_DIR}/hipify/gensrc # for rccl_bfloat16.h
     ${PROJECT_BINARY_DIR}/hipify/src # for graph/topo.h
     ${PROJECT_BINARY_DIR}/hipify/src/include/plugin # for recorder tests, nccl_tuner.h
+    ${PROJECT_BINARY_DIR}/hipify/src/transport/net_ib_cast # for net_ib_fault_inject.h, net_ib_cast_inspect.h
+    ${PROJECT_BINARY_DIR}/hipify/src/transport             # for net_ib_limits.h
     ${ROCM_PATH}/include
     ${ROCM_PATH}
   )