[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
Этот коммит содержится в:
isaki001
GitHub
+14
-10
@@ -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];
|
||||
|
||||
@@ -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);
|
||||
|
||||
+16
-5
@@ -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;
|
||||
|
||||
Ссылка в новой задаче
Block a user