Fix qpIndex selection in ncclIbIrecv for AINIC mode in net_ib_rocm

In AINIC mode, comm->base.qpIndex is intentionally not updated inside
the ncclIbIrecv recv-posting loop — it is deferred to ncclIbPostFifo
so that CTS messages are sent on the correct QPs. However, the loop
body was still using comm->base.qpIndex for QP selection, causing all
iterations to post receives on the same QP instead of distributing
them across all physical NICs in the merged device.

Introduce curQpIndex that reads from the local qpIndex variable (which
does advance each iteration) in AINIC mode, and from comm->base.qpIndex
in the standard path. This ensures round-robin QP selection works
correctly with both AINIC and non-AINIC configurations.

Author: amd-arozanov

projects/rccl/ext-src/rocm_netib.patch

@@ -665,7 +665,7 @@ index 9bfd8dcf..4d3f0a08 100644
    struct ncclIbRequest* req;
    NCCLCHECK(ncclIbGetRequest(&comm->base, &req));
    req->type = NCCL_NET_IB_REQ_RECV;
-@@ -2586,50 +2800,64 @@ ncclResult_t ncclIbIrecv(void* recvComm, int n, void** data, size_t* sizes, int*
+@@ -2586,50 +2800,65 @@ ncclResult_t ncclIbIrecv(void* recvComm, int n, void** data, size_t* sizes, int*
      req->devBases[i] = &comm->devs[i].base;
    }
 
@@ -697,7 +697,8 @@ index 9bfd8dcf..4d3f0a08 100644
 +    struct ibv_recv_wr* bad_wr;
 +    int qpIndex = comm->base.qpIndex;
 +    for (int i = 0; i < nqps; i++) {
-+      struct ncclIbQp* qp = comm->base.qps + comm->base.qpIndex;
++      int curQpIndex = rcclAinicRoce ? qpIndex : comm->base.qpIndex;
++      struct ncclIbQp* qp = comm->base.qps + curQpIndex;
 +      ncclIbAddEvent(req, qp->devIndex, &comm->devs[qp->devIndex].base);
  #ifdef NCCL_ENABLE_NET_PROFILING
 -    // Start a QP event for every request in the multirecv and every qp
@@ -718,7 +719,7 @@ index 9bfd8dcf..4d3f0a08 100644
 +      for (int r = 0; r < n; r++) {
 +        int nEventHandles = req->pInfo[r].nEventHandles;
 +        assert(nEventHandles < MAX_QPS_PER_REQ);
-+        req->pInfo[r].qpIndex[nEventHandles] = comm->base.qpIndex;
++        req->pInfo[r].qpIndex[nEventHandles] = curQpIndex;
 +        // Store info for profiler
 +        int64_t pluginId = NCCL_PROFILER_NET_TYPE_IB | NCCL_PROFILER_NET_IB_VER;
 +        req->pInfo[r].data.type = ncclProfileQp;
@@ -762,7 +763,7 @@ index 9bfd8dcf..4d3f0a08 100644
  }
 
  ncclResult_t ncclIbIflush(void* recvComm, int n, void** data, int* sizes, void** mhandles, void** request) {
-@@ -2698,6 +2926,8 @@ static int getReqQpIndex(struct ncclIbRequest* req, int request, int qpNumber) {
+@@ -2698,6 +2927,8 @@ static int getReqQpIndex(struct ncclIbRequest* req, int request, int qpNumber) {
  }
  #endif
 
@@ -771,7 +772,7 @@ index 9bfd8dcf..4d3f0a08 100644
  ncclResult_t ncclIbTest(void* request, int* done, int* sizes) {
    struct ncclIbRequest *r = (struct ncclIbRequest*)request;
    *done = 0;
-@@ -2731,13 +2961,18 @@ ncclResult_t ncclIbTest(void* request, int* done, int* sizes) {
+@@ -2731,13 +2962,18 @@ ncclResult_t ncclIbTest(void* request, int* done, int* sizes) {
 
      int totalWrDone = 0;
      int wrDone = 0;
@@ -792,7 +793,7 @@ index 9bfd8dcf..4d3f0a08 100644
          totalWrDone += wrDone;
          if (wrDone == 0) { TIME_CANCEL(3); } else { TIME_STOP(3); }
          if (wrDone == 0) continue;
-@@ -2889,7 +3124,7 @@ ncclResult_t rcclNetP2pPolicy(void* handle, int isP2p) {
+@@ -2889,7 +3125,7 @@ ncclResult_t rcclNetP2pPolicy(void* handle, int isP2p) {
  }
 
  ncclNet_t ncclNetIb = {

projects/rccl/src/transport/net_ib_rocm.cc

@@ -2819,14 +2819,15 @@ ncclResult_t rocmIbIrecv(void* recvComm, int n, void** data, size_t* sizes, int*
     struct ibv_recv_wr* bad_wr;
     int qpIndex = comm->base.qpIndex;
     for (int i = 0; i < nqps; i++) {
-      struct ncclIbQp* qp = comm->base.qps + comm->base.qpIndex;
+      int curQpIndex = rcclAinicRoce ? qpIndex : comm->base.qpIndex;
+      struct ncclIbQp* qp = comm->base.qps + curQpIndex;
       ncclIbAddEvent(req, qp->devIndex, &comm->devs[qp->devIndex].base);
 #ifdef NCCL_ENABLE_NET_PROFILING
       // Start a QP event for every request in the multirecv and every qp
       for (int r = 0; r < n; r++) {
         int nEventHandles = req->pInfo[r].nEventHandles;
         assert(nEventHandles < MAX_QPS_PER_REQ);
-        req->pInfo[r].qpIndex[nEventHandles] = comm->base.qpIndex;
+        req->pInfo[r].qpIndex[nEventHandles] = curQpIndex;
         // Store info for profiler
         int64_t pluginId = NCCL_PROFILER_NET_TYPE_IB | NCCL_PROFILER_NET_IB_VER;
         req->pInfo[r].data.type = ncclProfileQp;