From afd04a5117bc0ec61f1527b9ee283bea2a75a718 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Tue, 11 Mar 2025 09:32:09 -0700
Subject: [PATCH] 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: 4237caad6934872f441212a346007cc063981d0e]
---
 projects/rccl/src/graph/paths.cc      | 2 ++
 projects/rccl/test/common/EnvVars.cpp | 2 +-
 projects/rccl/test/common/TestBed.cpp | 2 +-
 3 files changed, 4 insertions(+), 2 deletions(-)

diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc
index e78026b2f2..a09df907dd 100644
--- a/projects/rccl/src/graph/paths.cc
+++ b/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
diff --git a/projects/rccl/test/common/EnvVars.cpp b/projects/rccl/test/common/EnvVars.cpp
index 3e0c73bb6e..15b697fabf 100644
--- a/projects/rccl/test/common/EnvVars.cpp
+++ b/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);
diff --git a/projects/rccl/test/common/TestBed.cpp b/projects/rccl/test/common/TestBed.cpp
index 612b2d68f1..7b9fbc88b7 100644
--- a/projects/rccl/test/common/TestBed.cpp
+++ b/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,