net-ib: fix fault injection API issues from review

- Validate qpIndex against comm->base.nqps (actual QP count) instead of
  NCCL_IB_MAX_QPS (128) in FaultSetQpDelay and FaultSetQpError, so
  arming a non-existent QP returns ncclInvalidArgument immediately.
- Reset fatalErrorCount in ncclIbCastFaultClear so the connection is
  truly recoverable after clearing fault state.
- Add phase 1 verification in FaultInjCastQpErrorClearRecovers: rank 1
  forwards sendRet + fatalCount to rank 0 via MPI, rank 0 asserts that
  the fault was actually observed before testing recovery in phase 2.

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

Author: i-kosarev

projects/rccl/src/transport/net_ib_cast.cc

@@ -2942,15 +2942,15 @@ static int IbCastQpSchedGetEffectiveTxNqps(struct ncclIbRequest* req, int *start
 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;
 }
@@ -2960,6 +2960,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/test/transport/NetIbMPI/FaultInjectTests.cpp

@@ -607,6 +607,10 @@ TEST_F(NetIbMPITest, FaultInjCastQpErrorClearRecovers) {
     MPI_Barrier(MPI_COMM_WORLD);
 
     // Trigger the fault (rank 0 posts recv, rank 1 posts send that will fail).
+    // rank 1 forwards its fault result to rank 0 so we can assert the fault
+    // was actually observed before testing recovery in phase 2.
+    static constexpr int kPhase1MpiTag = 9882;
+    FaultInjectResult p1 = {};
     if (rank == 0) {
         void*  bufs[1]    = {buf1};
         size_t sizes[1]   = {kMsgSize};
@@ -628,17 +632,32 @@ TEST_F(NetIbMPITest, FaultInjCastQpErrorClearRecovers) {
             if (sendRet != ncclSuccess || sendReq != nullptr) break;
             usleep(kPollIntervalUs);
         }
+        int fatalCount = 0;
         if (sendRet == ncclSuccess && sendReq != nullptr) {
             for (int poll = 0; poll < 200; poll++) {
                 int done = 0, sz = 0;
-                int fc = 0;
                 TestRequest(sendReq, &done, &sz);
-                ncclIbCastFaultGetFatalCount(sendComm1, &fc);
-                if (done || fc > 0) break;
+                ncclIbCastFaultGetFatalCount(sendComm1, &fatalCount);
+                if (done || fatalCount > 0) break;
                 usleep(kPollIntervalUs);
             }
+        } else {
+            ncclIbCastFaultGetFatalCount(sendComm1, &fatalCount);
         }
+        p1.sendRet    = static_cast<int>(sendRet);
+        p1.fatalCount = fatalCount;
         ncclIbCastFaultClear(sendComm1);
+        MPI_Send(&p1, sizeof(p1), MPI_BYTE, 0, kPhase1MpiTag, MPI_COMM_WORLD);
+    }
+
+    if (rank == 0) {
+        MPI_Recv(&p1, sizeof(p1), MPI_BYTE, 1, kPhase1MpiTag, MPI_COMM_WORLD,
+                 MPI_STATUS_IGNORE);
+        bool isendFailed = (p1.sendRet != static_cast<int>(ncclSuccess));
+        EXPECT_TRUE(isendFailed || p1.fatalCount > 0)
+            << "Phase 1: fault injection did not trigger — isend returned "
+            << p1.sendRet << ", fatalCount=" << p1.fatalCount
+            << "; recovery test is meaningless without a confirmed fault";
     }
 
     MPI_Barrier(MPI_COMM_WORLD);