From 44121db890387c409228c357751b54d30dfa5516 Mon Sep 17 00:00:00 2001 From: isaki001 <36317038+isaki001@users.noreply.github.com> Date: Fri, 15 Aug 2025 15:12:45 -0500 Subject: [PATCH] [TUNING] gfx950 16N tuning (#1835) * change gfx950 algo/proto selection for multinode allreduce, allgather, reduceScatter * gfx950 tuning: enable tuning for broadcast, allreduce starts LL128 earlier and switches to ring earlier, change LL128 start for allgather and reduceScatter * lower LL128 threshold * update reduceScatter LL128 min to match LL max for consistency * enable multinode PXN and increase chunksize for gfx950 * change LL128 start to 128KB, adjust ring-start according to node-count * disable code-path for fused-AR on LL128 for gfx950 * use LL128 starting from 1KB for multinode allgather on gfx950 * start LL128 earlier for multinode reduceScatter on gfx950 * start LL128 earlier for multinode broadcast on gfx950 * set multinode allreduce to start simple on 64MB for gfx950 * start LL128 from 1KB for multinode broadcast on gfx950 * setting multinode AR to use tree instead of ring at 16MB, 64MB, 128MB * set multinode broadcast to use LL for up to 256KB depending on node-count for gfx950 * adjust algo for 32MB multinode allreduce on gfx950 * make 32MB tree LL128 for multinode AR on gfx950 * make sure ring is not picked on 2N allreduce on small sizes --- src/graph/tuning.cc | 24 ++++++++++++++---------- src/include/rccl_common.h | 7 +++++-- src/rccl_wrap.cc | 21 ++++++++++++++++----- 3 files changed, 35 insertions(+), 17 deletions(-) diff --git a/src/graph/tuning.cc b/src/graph/tuning.cc index 682c220c79..be8f02730e 100644 --- a/src/graph/tuning.cc +++ b/src/graph/tuning.cc @@ -355,26 +355,30 @@ static struct tuningModel tuning_model_6 { { /* Tree (LL/LL128/Simple)*/ { 0.06, 0.06, 0.59 }, /* Ring (LL/LL128/Simple)*/ { 0.08, 0.08, 1.00 }, /* CollNetDirect (Simple)*/ { 0.00, 0.00, 1.00 }, /* CollNetChain (Simple)*/ { 0.00, 0.00, 1.00 }, /* NVLS */ { 0, 0, 0 }, /* NVLS Tree */ { 0, 0, 0 } }, }, - .treeCorrectionFactor = { + .treeCorrectionFactor = { /*16 32 64M 128M 1G*/ { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 1.0, 1.0, 1.0, 1.0, 0.6, 1.0, 0.9, 1.0, 1.0, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, }, - { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 1.0, 1.0, 1.0, 1.0, 0.6, 1.0, 0.9, 1.0, 1.0, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, }, - { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.7, 1.0, 1.0, 1.0, 1.0, 1.0, 0.7, 0.7, 0.8, 0.6, 0.6, 0.6, }, + { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 1.0, 1.0, 1.0, 1.0, 0.6, 1.0, 0.9, 1.0, 1.0, 1.0, 1.0, 1.0, 2.0, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, }, + { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.7, 1.0, 0.1, 0.1, 0.1, 0.9, 0.9, 0.7, 0.8, 0.6, 0.6, 0.6, }, }, .ringCorrectionFactor = { { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 0.2, 1.0, 0.4, 0.4, 0.1, 0.2, 0.2, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, }, - { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 0.2, 1.0, 0.4, 0.4, 0.1, 0.2, 0.2, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, }, - { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 0.8, 1.0, 1.0, 1.0, }, + { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 0.2, 1.0, 0.4, 0.4, 0.1, 0.2, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, }, + { 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 0.8, 0.8, 1.0, 1.0, }, }, // Follow order in RcclTunableColls .llProtoRanges = { /*ReduceScatter*/ - {/*LL (min/max/factor/thread_threshold)*/ {0, 65536, 1, 16}, /*LL64/128 (min/max/factor/thread_threshold)*/ {65536, 4194304, 1, 64}}, + {/*LL (min/max/factor/thread_threshold)*/ {0, 131071, 1, 16}, /*LL64/128 (min/max/factor/thread_threshold)*/ {131071, 4194304, 1, 64}}, /*AllGather*/ - {/*LL (min/max/factor/thread_threshold)*/ {0, 65535, 1, 16}, /*LL64/128 (min/max/factor/thread_threshold)*/ {65535, 8388608, 1, 64}}, + {/*LL (min/max/factor/thread_threshold)*/ {0, 7, 1, 16}, /*LL64/128 (min/max/factor/thread_threshold)*/ {7, 8388608, 1, 64}}, /*AllReduce*/ - {/*LL (min/max/factor/thread_threshold)*/ {0, 262144, 1, 0},/*LL64/128 (min/max/factor/thread_threshold)*/ {262144, 17660227, 3145728, 0}}, - }, + {/*LL (min/max/factor/thread_threshold)*/ {0, 131071, 1, 0},/*LL64/128 (min/max/factor/thread_threshold)*/ {131071, 17660227, 3145728, 0}}, + /*Reduce*/ + {/*LL (min/max/factor/thread_threshold)*/ {0, 16383, 1, 0},/*LL64/128 (min/max/factor/thread_threshold)*/ {16383, 16777216, 1, 0}}, + /*Broadcast*/ + {/*LL (min/max/factor/thread_threshold)*/ {0, 2048, 1, 0},/*LL64/128 (min/max/factor/thread_threshold)*/ {2048, 16777216, 1, 0}}, + }, }; static struct tuningModel rcclTuningModel[] = { @@ -824,8 +828,8 @@ ncclResult_t ncclTopoGetAlgoTime(struct ncclComm* comm, int coll, int algorithm, *time = -1.0; return ncclSuccess; } int logSize = log2i(nBytes>>6); - #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + if (algorithm == NCCL_ALGO_TREE) { if (logSize < 27) bw *= rcclTuningModel[comm->topo->tuning].treeCorrectionFactor[protocol][logSize]; else bw *= rcclTuningModel[comm->topo->tuning].treeCorrectionFactor[protocol][26]; diff --git a/src/include/rccl_common.h b/src/include/rccl_common.h index 914dc2c9ed..b1dd99114d 100644 --- a/src/include/rccl_common.h +++ b/src/include/rccl_common.h @@ -31,7 +31,8 @@ typedef enum RcclTunableColls { RCCL_AG_TUNABLE = 1, // all_gather index RCCL_AR_TUNABLE = 2, // all_reduce index RCCL_RE_TUNABLE = 3, // reduce index - RCCL_TUNABLE_COLLS = 4 // LL/LL64/LL128 tunable collectives count + RCCL_BR_TUNABLE = 4, // broadcast index + RCCL_TUNABLE_COLLS = 5 // LL/LL64/LL128 tunable collectives count } rcclTunableIndex_t; #define RCCL_LL_LIMITS_UNDEFINED 0 @@ -66,6 +67,8 @@ inline rcclTunableIndex_t rcclGetTunableIndex(ncclFunc_t const& func) { return RCCL_AR_TUNABLE; case ncclFuncReduce: return RCCL_RE_TUNABLE; + case ncclFuncBroadcast: + return RCCL_BR_TUNABLE; default: return RCCL_UNSUPPORTED_TUNABLE; // Invalid or unsupported function } @@ -76,7 +79,7 @@ inline size_t rcclGetSizePerRank(ncclFunc_t const& func, size_t const& nBytes, i // For AG, this is the send size per rank // For RS, this is the recv size per rank // For AR, this is the send/recv size per rank - return (func == ncclFuncReduceScatter || func == ncclFuncAllGather) ? nBytes / nRanks : nBytes; + return (func == ncclFuncReduceScatter || func == ncclFuncAllGather || func == ncclFuncBroadcast) ? nBytes / nRanks : nBytes; } void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info); void rcclUpdateThreadThreshold(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info, int& threadThreshold); diff --git a/src/rccl_wrap.cc b/src/rccl_wrap.cc index bf258a39ce..d1efc30bde 100644 --- a/src/rccl_wrap.cc +++ b/src/rccl_wrap.cc @@ -33,7 +33,7 @@ void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, s userProtocolInput = !protoStr ? 0 : 1; } - if(!userProtocolInput && comm->nNodes >= 2 && (info->func == ncclFuncReduceScatter || info->func == ncclFuncAllGather || info->func == ncclFuncAllReduce)) { + if(!userProtocolInput && comm->nNodes >= 2 && (info->func == ncclFuncReduceScatter || info->func == ncclFuncAllGather || info->func == ncclFuncAllReduce || info->func == ncclFuncBroadcast)) { auto tunableIndex = rcclGetTunableIndex(info->func); auto llMin = comm->minMaxLLRange[tunableIndex][NCCL_PROTO_LL][RCCL_PROTOCOL_MIN_IDX]; auto llMax = comm->minMaxLLRange[tunableIndex][NCCL_PROTO_LL][RCCL_PROTOCOL_MAX_IDX]; @@ -124,11 +124,14 @@ void rcclSetPxn(struct ncclComm* comm, int& rcclPxnDisable) { static int pxnDisable = RCCL_VALUE_UNSET; if(pxnDisable == RCCL_VALUE_UNSET) { const char *inputStr = getenv("NCCL_PXN_DISABLE"); - if(!IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || inputStr) { + const bool archGfx942 = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942"); + const bool archGfx950 = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950"); + if((!archGfx942 && !archGfx950) || inputStr) { rcclPxnDisable = pxnDisable = RCCL_VALUE_INVALID; return; } - pxnDisable = (comm->nRanks >= 64)? 0 : 1; + const int ranksThreshold = (archGfx942)? 64 : 32; + pxnDisable = (comm->nRanks >= ranksThreshold)? 0 : 1; INFO(NCCL_INIT, "RCCL PXN set as %s", !pxnDisable? "enabled" : "disabled"); } rcclPxnDisable = pxnDisable; @@ -138,11 +141,19 @@ void rcclSetP2pNetChunkSize(struct ncclComm* comm, int& rcclP2pNetChunkSize) { static int p2pNetChunkSize = RCCL_VALUE_UNSET; if(p2pNetChunkSize == RCCL_VALUE_UNSET) { const char *inputStr = getenv("NCCL_P2P_NET_CHUNKSIZE"); - if(!IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || inputStr) { + const bool archGfx942 = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942"); + const bool archGfx950 = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950"); + if((!archGfx942 && !archGfx950) || inputStr) { rcclP2pNetChunkSize = p2pNetChunkSize = RCCL_VALUE_INVALID; return; } - p2pNetChunkSize = (comm->nRanks >= 64)? (1 << 19) : (1 << 17); + + if(archGfx942) + p2pNetChunkSize = (comm->nRanks >= 64)? (1 << 19) : (1 << 17); + else if(archGfx950) + p2pNetChunkSize = (comm->nRanks >= 32) ? (1 << 19) : (comm->nRanks >= 16 ? (1 << 18) : (1 << 17)); + else + WARN("RCCL P2P attempt to set P2P net chunk size for unsupported arch: %s", comm->topo->nodes[GPU].nodes[0].gpu.gcn); INFO(NCCL_INIT, "RCCL P2P net chunk size default set to: %d", p2pNetChunkSize); } rcclP2pNetChunkSize = p2pNetChunkSize;