From eb0b1387b721ec78a7b9abd48031a8540c4486d2 Mon Sep 17 00:00:00 2001 From: Mustafa Abduljabbar Date: Thu, 30 Oct 2025 00:24:32 -0400 Subject: [PATCH] [Device] Adjust threadblock size for gfx950 to increase LL64/Simple performance for AR, RS and AG (#1978) * Add initial commit to increase tb size to 512 * Fix LL perf issue when subset of NCCL_MAX_NTHREADS is used Adding a constant to barrier_generic logic from using fallback logic when nthreads < NCCL_MAX_NTHREADS and nthreads == blockDim.X * Adjust nthreads for LL * Opt threads for reduce_scatter upper small range * Add macro for single node * Restrict MSCCL to 256 threads to prevent mem access fault * Support pre-MI350 compatibility * Partially refactor threadblock size override * Use const macros instead of numerals * opt out of unused function [ROCm/rccl commit: 12f51ba8bfa07c29a0c43a792446a962e1c8b9a1] --- projects/rccl/src/device/msccl_kernel_impl.h | 4 +- projects/rccl/src/device/primitives.h | 3 +- projects/rccl/src/device/prims_ll.h | 5 ++- projects/rccl/src/device/prims_ll128.h | 9 +++-- projects/rccl/src/device/prims_simple.h | 3 +- projects/rccl/src/enqueue.cc | 22 +++++++++-- projects/rccl/src/graph/tuning.cc | 27 +++++++------ projects/rccl/src/include/device.h | 24 +++++++++--- .../rccl/src/include/msccl/msccl_struct.h | 4 +- projects/rccl/src/include/rccl_common.h | 9 +++++ projects/rccl/src/init.cc | 10 ++++- projects/rccl/src/misc/msccl/msccl_setup.cc | 2 +- projects/rccl/src/rccl_wrap.cc | 38 +++++++++++++++++-- 13 files changed, 120 insertions(+), 40 deletions(-) diff --git a/projects/rccl/src/device/msccl_kernel_impl.h b/projects/rccl/src/device/msccl_kernel_impl.h index d2419b4f15..87ba510ba4 100644 --- a/projects/rccl/src/device/msccl_kernel_impl.h +++ b/projects/rccl/src/device/msccl_kernel_impl.h @@ -28,7 +28,7 @@ extern __shared__ struct mscclShmemData mscclShmem; inline __device__ static void barrier(int nthreads) { #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) - assert(nthreads == NCCL_MAX_NTHREADS); + assert(nthreads == MSCCL_MAX_NTHREADS); #ifdef __GFX12__ __asm__ __volatile__("s_waitcnt vmcnt(0) lgkmcnt(0)\ns_barrier_signal -1\ns_barrier_wait -1"); #else @@ -96,7 +96,7 @@ __device__ __forceinline__ void mscclRunInterpreter( struct ncclDevComm* comm, struct mscclAlgo* algo, struct mscclWork* work) { const int tid = threadIdx.x; const int bid = blockIdx.x; - const int nthreads = NCCL_MAX_NTHREADS; + const int nthreads = MSCCL_MAX_NTHREADS; #if defined(ENABLE_NPKIT) uint64_t timestamp_entry = 0; diff --git a/projects/rccl/src/device/primitives.h b/projects/rccl/src/device/primitives.h index 28ecc0d36c..b09c7b5031 100644 --- a/projects/rccl/src/device/primitives.h +++ b/projects/rccl/src/device/primitives.h @@ -16,9 +16,8 @@ #include "common.h" #define NCCL_SPINS_BEFORE_CHECK_ABORT 10000 - #define barrier_generic(__THREAD_FENCE, NWORKERS, BARRIER_NEXT, BARRIERS_PTR) do { \ - if (nthreads == NCCL_MAX_NTHREADS) { \ + if (nthreads == threadsPerBlock) { \ __THREAD_FENCE; __builtin_amdgcn_s_barrier(); \ } else { \ /**const int w = threadIdx.x/WARP_SIZE //unused variable - compiler warning**/;\ diff --git a/projects/rccl/src/device/prims_ll.h b/projects/rccl/src/device/prims_ll.h index 1d6ba34822..0730ee2694 100644 --- a/projects/rccl/src/device/prims_ll.h +++ b/projects/rccl/src/device/prims_ll.h @@ -25,6 +25,7 @@ class PrimitivesworkBytes; size_t batchBytes = plan->nWorkBatches*sizeof(struct ncclDevWorkBatch); +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#else plan->threadPerBlock = std::max(plan->threadPerBlock, 256 /*NCCL_MIN_NTHREADS*/); - +#endif // If we can fit everything into the kernel args we do so. if (sizeof(ncclDevKernelArgs) + batchBytes + workBytes <= comm->workArgsBytes) { plan->workStorageType = ncclDevWorkStorageTypeArgs; @@ -866,7 +869,11 @@ static ncclResult_t scheduleCollTasksToPlan( plan->channelMask.masks[maskIdx] |= (1ull<channelMask.masks[channelId/64] |= (2ull<channelHi) - (1ull<channelLo); +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + plan->threadPerBlock = task->nWarps * comm->WarpSize; +#else plan->threadPerBlock = std::max(plan->threadPerBlock, 192 /* 3*WARP_SIZE */); +#endif if (!plan->kernelSpecialized) { plan->kernelFn = ncclKerns[ncclGetKernelIndex(comm)].kernelFn; plan->kernelSpecialized = ncclKerns[ncclGetKernelIndex(comm)].specialized; @@ -1200,8 +1207,11 @@ static ncclResult_t scheduleP2pTasksToPlan( ) { int nRanks = comm->nRanks; struct ncclKernelPlanner::Peer* peers = comm->planner.peers; - +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + plan->threadPerBlock = std::max(plan->threadPerBlock, RCCL_P2P_MAX_NTHREADS); +#else plan->threadPerBlock = std::max(plan->threadPerBlock, NCCL_MAX_NTHREADS); +#endif if (!plan->kernelSpecialized) { plan->kernelFn = ncclKerns[ncclGetKernelIndex(comm)].kernelFn; plan->kernelSpecialized = ncclKerns[ncclGetKernelIndex(comm)].specialized; @@ -2133,8 +2143,12 @@ static ncclResult_t topoGetAlgoInfo( } else { info->nMaxChannels = nc; } +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#else if (info->algorithm == NCCL_ALGO_TREE) nt = NCCL_MAX_NTHREADS; // Tree now uses all threads always. - if (info->algorithm == NCCL_ALGO_PAT) nt = NCCL_MAX_NTHREADS; + if (info->algorithm == NCCL_ALGO_PAT) nt = NCCL_MAX_NTHREADS; +#endif + rcclOptThreadBlockSize(comm, info, nBytes, nt); info->nWarps = nt/comm->WarpSize; rcclOverrideAlgorithm(ncclAlgoStr, table, info); rcclOverrideProtocol(ncclProtoStr, table, info); diff --git a/projects/rccl/src/graph/tuning.cc b/projects/rccl/src/graph/tuning.cc index dba12f3b26..fe6ac8ec78 100644 --- a/projects/rccl/src/graph/tuning.cc +++ b/projects/rccl/src/graph/tuning.cc @@ -362,7 +362,7 @@ static struct tuningModel tuning_model_5 { }, .channelThresholds = {{{CHAN_THRESHOLDS_UNDEFINED}}}, - + }; static struct tuningModel tuning_model_6 { @@ -405,14 +405,14 @@ static struct tuningModel tuning_model_6 { {/*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}}, - }, - + }, + .channelThresholds = { // For each collective, define minMax per-rank size threshold for 32,40,48,56,64 channels /*ReduceScatter*/ {{512, 1024, 2},{1024, 2048, 4},{2048, 4096, 8},{4096, 65536, 16}, {65536, 262144, 32}, {262144, 524288, 40}, {1,1, 48}, {524288, 1048576, 56}, {1048576, 268435457, 64}}, /*AllGather*/ {{2048, 4096, 2},{4096, 8192, 4},{8192, 16384, 8},{16384, 262144, 16},{262144, 524288, 32}, {524288, 1048576, 40}, {1,1, 48}, {1048576, 4194304, 56}, {4194304, 268435457, 64}}, /*AllReduce*/ {{0,0,0},{0,0,0},{0,0,0},{0,0,0},{0,0,0}, {0,0,0}, {0,0,0}, {0,0,0}, {0,0,0}}, - }, + }, }; static struct tuningModel rcclTuningModel[] = { @@ -484,18 +484,23 @@ static float getNetOverhead(struct ncclComm* comm) { NCCL_PARAM(Ll128C2c, "LL128_C2C", 1); ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph** graphs) { - int simpleDefaultThreads = (graphs[NCCL_ALGO_RING]->bwIntra*graphs[NCCL_ALGO_RING]->nChannels <= PCI_BW) ? 256 : NCCL_SIMPLE_MAX_NTHREADS; - comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_SIMPLE] = #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) - getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, NCCL_MAX_NTHREADS, simpleDefaultThreads, comm->WarpSize); + static int rcclMaxThreads[NCCL_NUM_PROTOCOLS] = {0}; + if (rcclMaxThreads[NCCL_PROTO_SIMPLE] == 0) rcclGetMaxNthreads(comm, rcclMaxThreads); + static int maxNthreads = rcclMaxThreads[NCCL_PROTO_SIMPLE]; + static int maxLL128Nthreads = rcclMaxThreads[NCCL_PROTO_LL128]; + static int maxLLThreads = rcclMaxThreads[NCCL_PROTO_LL]; + int simpleDefaultThreads = (graphs[NCCL_ALGO_RING]->bwIntra*graphs[NCCL_ALGO_RING]->nChannels <= PCI_BW) ? 256 : maxNthreads; + comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_SIMPLE] = getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, maxNthreads, simpleDefaultThreads, comm->WarpSize); comm->maxThreads[NCCL_ALGO_TREE][NCCL_PROTO_SIMPLE] = comm->maxThreads[NCCL_ALGO_COLLNET_DIRECT][NCCL_PROTO_SIMPLE] = - getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, NCCL_MAX_NTHREADS, NCCL_MAX_NTHREADS, comm->WarpSize); + getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, maxNthreads, maxNthreads, comm->WarpSize); comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_LL] = comm->maxThreads[NCCL_ALGO_TREE][NCCL_PROTO_LL] = comm->maxThreads[NCCL_ALGO_COLLNET_DIRECT][NCCL_PROTO_LL] = - getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, NCCL_MAX_NTHREADS, NCCL_MAX_NTHREADS, comm->WarpSize); + getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, maxNthreads, maxLLThreads, comm->WarpSize); comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_LL128] = comm->maxThreads[NCCL_ALGO_TREE][NCCL_PROTO_LL128] = - getNthreads("NCCL_LL128_NTHREADS", ncclParamLl128Nthreads(), 4*comm->WarpSize, NCCL_LL128_MAX_NTHREADS, NCCL_LL128_MAX_NTHREADS, comm->WarpSize); + getNthreads("NCCL_LL128_NTHREADS", ncclParamLl128Nthreads(), 4*comm->WarpSize, maxLL128Nthreads, maxLL128Nthreads, comm->WarpSize); #else - getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 2*WARP_SIZE, NCCL_SIMPLE_MAX_NTHREADS, simpleDefaultThreads); + int simpleDefaultThreads = (graphs[NCCL_ALGO_RING]->bwIntra*graphs[NCCL_ALGO_RING]->nChannels <= PCI_BW) ? 256 : NCCL_MAX_NTHREADS; + comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_SIMPLE] = getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 2*WARP_SIZE, NCCL_SIMPLE_MAX_NTHREADS, simpleDefaultThreads); comm->maxThreads[NCCL_ALGO_TREE][NCCL_PROTO_SIMPLE] = getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 2*WARP_SIZE, NCCL_SIMPLE_MAX_NTHREADS, NCCL_SIMPLE_MAX_NTHREADS); comm->maxThreads[NCCL_ALGO_COLLNET_DIRECT][NCCL_PROTO_SIMPLE] = diff --git a/projects/rccl/src/include/device.h b/projects/rccl/src/include/device.h index 3e0cb12f3c..8b420d7f1b 100644 --- a/projects/rccl/src/include/device.h +++ b/projects/rccl/src/include/device.h @@ -101,19 +101,28 @@ union ncclLLFifoLine { #else #define WARP_SIZE 32 #endif + #if defined (__gfx950__) + #define NCCL_MAX_NTHREADS 512 + #else + #define NCCL_MAX_NTHREADS 256 + #endif + // Number of named barriers supported by CUDA + #define NCCL_MAX_GROUPS (NCCL_MAX_NTHREADS/WARP_SIZE) #else - /* IMPORTANT: - * WARP_SIZE should NEVER be referenced by host code in RCCL. It is defined here + /* IMPORTANT Note ragarding WARP_SIZE, NCCL_MAX_NTHREADS and NCCL_MAX_GROUPS: + * These should NEVER be referenced by host code in RCCL. It is defined here * solely as a workaround to allow RCCL to compile, since the host still compiles __device__ functions, - * and WARP_SIZE needs to be defined. These __device__ functions will not be called from the host. + * and they need to be defined. These __device__ functions will not be called from the host. * The host warp size is handled in src/enqueue.cc by calling hipDeviceGetAttributes(). */ #define WARP_SIZE 32 + #define NCCL_MAX_NTHREADS 256 + // Number of named barriers supported by CUDA + #define NCCL_MAX_GROUPS (NCCL_MAX_NTHREADS/WARP_SIZE) #endif #define MAXCHANNELS 128 #define CHANNEL_LIMIT 16 #define NCCL_MAX_LOCAL_RANKS 72 -#define NCCL_MAX_NTHREADS 256 #define NCCL_MIN_NTHREADS (4*WARP_SIZE) #define NCCL_SIMPLE_MAX_NTHREADS NCCL_MAX_NTHREADS #define NCCL_SIMPLE_EXTRA_GROUP_IF_NTHREADS_GE (3*WARP_SIZE) @@ -145,8 +154,7 @@ static_assert(NCCL_LL_CLEAN_MASK % NCCL_STEPS == 0, "Invalid NCCL_LL_CLEAN_MASK #define NCCL_DIRECT_NIC 0x04 #define NCCL_NVLS_MIN_POLL 0x80 -// Number of named barriers supported by CUDA -#define NCCL_MAX_GROUPS (NCCL_MAX_NTHREADS/WARP_SIZE) + #define NCCL_REGULAR_BUFFER 0x00 #define NCCL_IPC_REG_BUFFER 0x01 @@ -679,10 +687,14 @@ __device__ constexpr int ncclShmemScratchWarpSize(int cudaArch = NCCL_CUDA_ARCH) ) + 15) & -16; // pad to 16 bytes } +// RCCL has its own varient of ncclShmemDynamicSize and ncclShmemScratchWarpSize +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#else // The amount of dynamic shmem per block __device__ constexpr int ncclShmemDynamicSize(int cudaArch = NCCL_CUDA_ARCH) { return cudaArch < 700 ? 0 : ncclShmemScratchWarpSize(cudaArch)*(NCCL_MAX_NTHREADS/WARP_SIZE); } +#endif // Host-side table of kernel function pointers. extern int const ncclDevKernelCount; diff --git a/projects/rccl/src/include/msccl/msccl_struct.h b/projects/rccl/src/include/msccl/msccl_struct.h index 351a557b11..6b4dbd06a7 100644 --- a/projects/rccl/src/include/msccl/msccl_struct.h +++ b/projects/rccl/src/include/msccl/msccl_struct.h @@ -36,6 +36,8 @@ #define MSCCL_LOCAL_COPY 6 #define MSCCL_REDUCE 7 +#define MSCCL_MAX_NTHREADS 256 + struct mscclTransmission { int16_t dependencePointer; // index to the first dependence int16_t numDependencies; // dependencePointer+numDependencies indicate the last dependence @@ -174,7 +176,7 @@ enum mscclCaptureStatus { struct mscclProxyArg { struct mscclAlgo* hostAlgo; ncclComm_t comm; - mscclProxyArg(struct mscclAlgo* hostAlgo, ncclComm_t comm) + mscclProxyArg(struct mscclAlgo* hostAlgo, ncclComm_t comm) : hostAlgo(hostAlgo), comm(comm) {} }; diff --git a/projects/rccl/src/include/rccl_common.h b/projects/rccl/src/include/rccl_common.h index 66c8733bae..7eb4dd6f90 100644 --- a/projects/rccl/src/include/rccl_common.h +++ b/projects/rccl/src/include/rccl_common.h @@ -45,6 +45,12 @@ typedef enum RcclTunableColls { #define RCCL_PROTOCOL_FACTOR_IDX 2 #define RCCL_PROTOCOL_THREAD_THRESHOLD_IDX 3 +#define RCCL_SINGLE_NODE_MAX_NTHREADS 256 +#define RCCL_GFX950_MAX_NTHREADS 512 // for Simple and LL64/LL128 gfx950 +#define RCCL_DEFAULT_MAX_NTHREADS 256 // for Simple and LL64/LL128 other archs +#define RCCL_LL_MAX_NTHREADS 256 +#define RCCL_P2P_MAX_NTHREADS 256 + typedef enum { RCCL_VALUE_UNSET = -2, RCCL_VALUE_INVALID = -1 @@ -98,6 +104,9 @@ ncclResult_t rcclOverrideAlgorithm(const char* ncclAlgoStr[], float table[][NCCL 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); void rcclSetPipelining(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info); +void rcclGetMaxNthreads(struct ncclComm* comm, int maxNthreads[]); +void rcclOptThreadBlockSize(struct ncclComm* comm, struct ncclTaskColl* info, size_t nBytes, int& nThreads); +void rcclSetDefaultBuffSizes(struct ncclComm* comm, int defaultBuffSizes[]); NCCL_API(ncclResult_t, rcclGetAlgoInfo, struct ncclComm* comm, ncclFunc_t coll, uint64_t count, ncclDataType_t dataType, int collNetSupport, int nvlsSupport, int numPipeOps, int* algo, int* protocol, int* maxChannels); NCCL_API(ncclResult_t, rcclGetAlgoName, int algo, const char** algoName); NCCL_API(ncclResult_t, rcclGetProtocolName, int protocol, const char** algoName); diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index b31d87029f..4bfcb85cdc 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -1007,10 +1007,12 @@ static ncclResult_t setupChannel(struct ncclComm* comm, int channelId, int rank, } return ncclSuccess; } - +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#else #define DEFAULT_LL_BUFFSIZE (NCCL_LL_LINES_PER_THREAD*NCCL_LL_MAX_NTHREADS*NCCL_STEPS*sizeof(union ncclLLFifoLine)) #define DEFAULT_LL128_BUFFSIZE (NCCL_LL128_ELEMS_PER_THREAD*NCCL_LL128_MAX_NTHREADS*NCCL_STEPS*sizeof(uint64_t)) #define DEFAULT_BUFFSIZE (1 << 22) /* 4MiB */ +#endif NCCL_PARAM(BuffSize, "BUFFSIZE", -2); NCCL_PARAM(LlBuffSize, "LL_BUFFSIZE", -2); NCCL_PARAM(Ll128BuffSize, "LL128_BUFFSIZE", -2); @@ -1021,8 +1023,12 @@ NCCL_PARAM(P2pNvlChunkSize, "P2P_NVL_CHUNKSIZE", (1 << 19)); /* 512 kB */ static ncclResult_t computeBuffSizes(struct ncclComm* comm) { int64_t envs[NCCL_NUM_PROTOCOLS] = { ncclParamLlBuffSize(), ncclParamLl128BuffSize(), ncclParamBuffSize() }; +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + int defaults[NCCL_NUM_PROTOCOLS]; + rcclSetDefaultBuffSizes(comm, defaults); +#else int defaults[NCCL_NUM_PROTOCOLS] = { DEFAULT_LL_BUFFSIZE, DEFAULT_LL128_BUFFSIZE, DEFAULT_BUFFSIZE }; - +#endif for (int p=0; pbuffSizes[p] = envs[p] != -2 ? envs[p] : defaults[p]; } diff --git a/projects/rccl/src/misc/msccl/msccl_setup.cc b/projects/rccl/src/misc/msccl/msccl_setup.cc index 1cae215c41..ec03dff009 100644 --- a/projects/rccl/src/misc/msccl/msccl_setup.cc +++ b/projects/rccl/src/misc/msccl/msccl_setup.cc @@ -420,7 +420,7 @@ ncclResult_t mscclSetupKernel(const void* sendBuff, void* recvBuff, size_t count uint32_t numBlocks = (uint32_t)hostAlgo->nBlocks; dim3 grid = {numBlocks, 1, 1}; - dim3 block = {NCCL_MAX_NTHREADS, 1, 1}; + dim3 block = {MSCCL_MAX_NTHREADS, 1, 1}; ncclDevRedOpFull opFull = {}; NCCLCHECK(hostToDevRedOp(&opFull, op, dataType, comm)); diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index 36eb728441..73bb816ac6 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -130,7 +130,7 @@ ncclResult_t rcclOverrideChannels(struct ncclComm* comm, ncclFunc_t coll, size_t int maxNChannels = std::max(comm->nChannels, static_cast(ncclParamMaxNchannels())); size_t bytesPerRank = divUp(nBytes, comm->nRanks); - for(int channelCountIndex = 0; channelCountIndex < RCCL_CHANNELS_TUNABLE_ENTRIES; ++channelCountIndex){ + for(int channelCountIndex = 0; channelCountIndex < RCCL_CHANNELS_TUNABLE_ENTRIES; ++channelCountIndex){ size_t minByteThreshold = comm->minMaxChannelThresholds[tunableIndex][channelCountIndex][0]; size_t maxByteThreshold = comm->minMaxChannelThresholds[tunableIndex][channelCountIndex][1]; INFO(NCCL_TUNING, "nBytes:%lu bytesPerRank:%lu minByteThreshold:%lu maxByteThreshold:%lu NCCL_MIN_NCHANNELS:%i or NCCL_MAX_NCHANNELS:%i minCTAs:%i maxCTAs:%i", nBytes, bytesPerRank, minByteThreshold, maxByteThreshold, minNChannels, maxNChannels, minCTAs, maxCTAs); @@ -138,7 +138,7 @@ ncclResult_t rcclOverrideChannels(struct ncclComm* comm, ncclFunc_t coll, size_t INFO(NCCL_TUNING, "RCCL tuning model does not define threshold for coll:%i and nbytes:%lu", coll, nBytes); break; // Skip undefined thresholds } - + if(bytesPerRank > minByteThreshold && bytesPerRank <= maxByteThreshold){ int channelCount = comm->minMaxChannelThresholds[tunableIndex][channelCountIndex][2]; @@ -353,13 +353,13 @@ bool rcclUseAllGatherDirect(struct ncclComm* comm, size_t& msgSize) { threshold = comm->nNodes * 2097152; } } else if (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942")) { - threshold = 4194304; + threshold = 4194304; } comm->enableCustColl = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942"); int rankMultiple = comm->nRanks % 8; - + //return (comm->enableCustColl && (comm->nNodes > 1) && (msgSize <= threshold) && (threshold != -1)) return (comm->enableCustColl && (msgSize <= threshold) && (threshold != -1) && !rankMultiple) ; @@ -408,6 +408,36 @@ void rcclSetP2pNetChunkSize(struct ncclComm* comm, int& rcclP2pNetChunkSize) { rcclP2pNetChunkSize = p2pNetChunkSize; } +void rcclGetMaxNthreads(struct ncclComm* comm, int maxNthreads[]) { + if (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) { + maxNthreads[NCCL_PROTO_SIMPLE] = maxNthreads[NCCL_PROTO_LL128] = RCCL_GFX950_MAX_NTHREADS; + } else { + maxNthreads[NCCL_PROTO_SIMPLE] = maxNthreads[NCCL_PROTO_LL128] = RCCL_DEFAULT_MAX_NTHREADS; + } + maxNthreads[NCCL_PROTO_LL] = RCCL_LL_MAX_NTHREADS; +} + +void rcclOptThreadBlockSize(struct ncclComm* comm, struct ncclTaskColl* info, size_t nBytes, int& nThreads) { + static int maxNthreads[NCCL_NUM_PROTOCOLS] = {0}; + if (maxNthreads[NCCL_PROTO_SIMPLE] == 0) rcclGetMaxNthreads(comm, maxNthreads); + if (info->algorithm == NCCL_ALGO_TREE) nThreads = maxNthreads[NCCL_PROTO_SIMPLE]; // Tree now uses all threads always. + if (info->algorithm == NCCL_ALGO_PAT) nThreads = maxNthreads[NCCL_PROTO_SIMPLE]; + if (comm->nNodes == 1) nThreads = RCCL_SINGLE_NODE_MAX_NTHREADS; // For single node, we use half the number of threads for perf reasons. + // The following should be already set correctly by getNthreads + // but need to override the changes for TREE and PAT in the previous lines + if (info->protocol == NCCL_PROTO_LL) nThreads = maxNthreads[NCCL_PROTO_LL]; + // ReduceScatter small count optimization + if (info->func == ncclFuncReduceScatter && divUp(nBytes, comm->nRanks) <= 524288) nThreads = maxNthreads[NCCL_PROTO_LL]; +} + +void rcclSetDefaultBuffSizes(struct ncclComm* comm, int defaultBuffSizes[]) { + static int maxNthreads[NCCL_NUM_PROTOCOLS] = {0}; + if (maxNthreads[NCCL_PROTO_SIMPLE] == 0) rcclGetMaxNthreads(comm, maxNthreads); + defaultBuffSizes[NCCL_PROTO_LL] = NCCL_LL_LINES_PER_THREAD*maxNthreads[NCCL_PROTO_LL]*NCCL_STEPS*sizeof(union ncclLLFifoLine); + defaultBuffSizes[NCCL_PROTO_LL128] = NCCL_LL128_ELEMS_PER_THREAD*maxNthreads[NCCL_PROTO_LL128]*NCCL_STEPS*sizeof(uint64_t); + defaultBuffSizes[NCCL_PROTO_SIMPLE] = (1 << 22); /* 4MiB */ +} + ncclResult_t rcclFuncMaxSendRecvCount(ncclFunc_t func, int nRanks, size_t count, size_t& maxCount) { RCCL_STATIC_EXPOSE_CHECK(); maxCount = ncclFuncMaxSendRecvCount(func, nRanks, count);