diff --git a/src/enqueue.cc b/src/enqueue.cc index c087aef5e1..89bab1f18d 100644 --- a/src/enqueue.cc +++ b/src/enqueue.cc @@ -893,7 +893,7 @@ static ncclResult_t addP2pToPlan( connIndex[dir] = NCCL_CONN_IDX_P2P_NET; } } - + if (!selfSend) { for (int part=0; part < nChannelsMax; part++) { int channelId = ncclP2pChannelForPart(comm->p2pnChannels, base, part, nChannelsMax, comm->nNodes); @@ -1679,6 +1679,14 @@ static ncclResult_t updateCollCostTable( return ncclSuccess; } +// The following parameters are based on the observation that LL and LL128 +// performance is determined by how much data goes out of a GPU +// TODO: these numbers can be part of topo detection. +RCCL_PARAM(RsLLMinSizePerRank, "REDUCE_SCATTER_LL_MIN_SIZE_PER_RANK", 0); +RCCL_PARAM(RsLLMaxSizePerRank, "REDUCE_SCATTER_LL_MAX_SIZE_PER_RANK", 655360); +RCCL_PARAM(RsLL128MinSizePerRank, "REDUCE_SCATTER_LL128_MIN_SIZE_PER_RANK", 131072); +RCCL_PARAM(RsLL128MaxSizePerRank, "REDUCE_SCATTER_LL128_MAX_SIZE_PER_RANK", 3211264); + static ncclResult_t topoGetAlgoInfo( struct ncclComm* comm, struct ncclTaskColl* info, size_t nBytes, float** collCostTable, int backupAlgo, int backupProto, float backupTime, ncclSimInfo_t* simInfo @@ -1712,6 +1720,32 @@ static ncclResult_t topoGetAlgoInfo( info->protocol = backupProto; time = backupTime; } +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + // Honor user input for protocol choice + static int userProtocolInput = -2; + if (userProtocolInput == -2) { + const char *protoStr = getenv("NCCL_PROTO"); + userProtocolInput = !protoStr ? 0 : 1; + } + if(!userProtocolInput && comm->nNodes >= 2 && info->func == ncclFuncReduceScatter && IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942")) { + // Keep it simple unless otherwise required + info->protocol = NCCL_PROTO_SIMPLE; + // Normalize the comparison to sizePerRank as this is essentially what matters in determining protocol choice + size_t sizePerRank = nBytes / comm->nRanks; + + if(sizePerRank <= rcclParamRsLLMaxSizePerRank() && sizePerRank >= rcclParamRsLLMinSizePerRank()) { + info->protocol = NCCL_PROTO_LL; + } +#if defined(ENABLE_LL128) + // LL128 RS performance is better than LL when enabled, so the next condition overrides the previous LL choice + if(comm->topo->ll128Enabled) { + if(sizePerRank <= rcclParamRsLL128MaxSizePerRank() && sizePerRank >= rcclParamRsLL128MinSizePerRank()) { + info->protocol = NCCL_PROTO_LL128; + } + } +#endif + } +#endif if (comm->rank == 0) INFO(NCCL_TUNING, "%ld Bytes -> Algo %d proto %d time %f", nBytes, info->algorithm, info->protocol, time); if (simInfo) simInfo->estimatedTime = time; TRACE(NCCL_COLL, "%ld Bytes -> Algo %d proto %d time %f", nBytes, info->algorithm, info->protocol, time);