Limit P2P channels per peer to not exceeding max channels (#1594)

* Limit P2P channels per peer to not exceeding max channels

* [UT] test single GPU cases for all collectives

* [UT] fix out of range root value

[ROCm/rccl commit: 4237caad69]

projects/rccl/src/graph/paths.cc

@@ -908,6 +908,8 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) {
     // Doubling P2P channels per peer on single node
     if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950"))) comm->p2pnChannelsPerPeer *= 2;
     comm->p2pnChannels = std::min(pow2Up(comm->p2pnChannels), 4*CHANNEL_LIMIT);
+    // p2pnChannelsPerPeer cannot be greater than MAXCHANNELS
+    comm->p2pnChannelsPerPeer = std::min(comm->p2pnChannelsPerPeer, MAXCHANNELS);
   }
 
   // Init channels that weren't used so far

projects/rccl/test/common/EnvVars.cpp

@@ -204,7 +204,7 @@ namespace RcclUnitTesting
     getArchInfo(&isGfx12, "gfx12");
 
     showNames      = GetEnvVar("UT_SHOW_NAMES"  , 1);
-    minGpus        = GetEnvVar("UT_MIN_GPUS"    , 2);
+    minGpus        = GetEnvVar("UT_MIN_GPUS"    , 1);
     maxGpus        = GetEnvVar("UT_MAX_GPUS"    , numDetectedGpus);
     processMask    = GetEnvVar("UT_PROCESS_MASK", UT_SINGLE_PROCESS | UT_MULTI_PROCESS);
     verbose        = GetEnvVar("UT_VERBOSE"     , 0);

projects/rccl/test/common/TestBed.cpp

@@ -695,7 +695,7 @@ namespace RcclUnitTesting
                                                     &numInputElements,
                                                     &numOutputElements);
           optionalArgs.redOp = redOps[rdIdx];
-          optionalArgs.root = roots[rtIdx];
+          optionalArgs.root = roots[rtIdx] % this->numActiveRanks;
           this->SetCollectiveArgs(funcTypes[ftIdx],
                                   dataTypes[dtIdx],
                                   numInputElements,