From 2447161fb48369a63786ca35cb4f5745fbf3d934 Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Wed, 22 May 2024 16:27:59 -0500 Subject: [PATCH] graph: restrict MAXCHANNELS for certain platforms [ROCm/rccl commit: 9f654f6cf557c400807bf4f8c2bbdd6c8355a4c3] --- projects/rccl/src/graph/connect.cc | 4 ++-- projects/rccl/src/include/device.h | 1 + 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/rccl/src/graph/connect.cc b/projects/rccl/src/graph/connect.cc index a87fc92749..55b3a6491e 100644 --- a/projects/rccl/src/graph/connect.cc +++ b/projects/rccl/src/graph/connect.cc @@ -282,7 +282,7 @@ static ncclResult_t setTreeDown(struct ncclTree* tree, int* indexes, int d) { static ncclResult_t connectTrees(struct ncclComm* comm, int* treeToParent, int* treeToChild0, int* treeToChild1, int* treePatterns) { - const int channelLimit = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ? MAXCHANNELS/2 : 16; + const int channelLimit = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ? 2*CHANNEL_LIMIT : CHANNEL_LIMIT; const int nChannels = (comm->nChannels > channelLimit) ? comm->nChannels / 2 : comm->nChannels; const int nNodes = comm->nNodes, node = comm->node; @@ -625,7 +625,7 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa NCCLCHECK(connectTrees(comm, treeToParent, treeToChild0, treeToChild1, treePatterns)); // Only use full MAXCHANNELS for gfx94x - int maxChannels = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ? MAXCHANNELS : (MAXCHANNELS/2); + int maxChannels = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ? MAXCHANNELS : 2*CHANNEL_LIMIT; // Duplicate ringPrev/ringNext for ncclBuildRing if (nChannels <= maxChannels/2) memcpy(ringPrev+nChannels*nranks, ringPrev, nChannels*nranks*sizeof(int)); diff --git a/projects/rccl/src/include/device.h b/projects/rccl/src/include/device.h index c4e88a1595..bd5957fb7e 100644 --- a/projects/rccl/src/include/device.h +++ b/projects/rccl/src/include/device.h @@ -62,6 +62,7 @@ union ncclLLFifoLine { #define WARP_SIZE warpSize #define MAXCHANNELS 256 +#define CHANNEL_LIMIT 16 #define NCCL_MAX_NTHREADS 256 #define NCCL_SIMPLE_MAX_NTHREADS NCCL_MAX_NTHREADS