From f773748b464ea76930d3aa4cd24f270f6c955cb8 Mon Sep 17 00:00:00 2001 From: David Addison Date: Mon, 13 Sep 2021 14:43:22 -0700 Subject: [PATCH] Resync with NCCL 2.11 New operator: mulsum New test: gather --- src/Makefile | 2 +- src/common.cu | 201 +++++++++++++++++++++++++++++++++++--------------- src/common.h | 9 ++- src/gather.cu | 131 ++++++++++++++++++++++++++++++++ 4 files changed, 279 insertions(+), 64 deletions(-) create mode 100644 src/gather.cu diff --git a/src/Makefile b/src/Makefile index 26e653e7d6..c8491ea537 100644 --- a/src/Makefile +++ b/src/Makefile @@ -70,7 +70,7 @@ NVLDFLAGS += $(LIBRARIES:%=-l%) DST_DIR := $(BUILDDIR) SRC_FILES := $(wildcard *.cu) OBJ_FILES := $(SRC_FILES:%.cu=${DST_DIR}/%.o) -BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall scatter sendrecv hypercube +BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall scatter gather sendrecv hypercube BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf) build: ${BIN_FILES} diff --git a/src/common.cu b/src/common.cu index 6a26c6c4e8..05f814d923 100644 --- a/src/common.cu +++ b/src/common.cu @@ -14,37 +14,37 @@ int test_ncclVersion = 0; // init'd with ncclGetVersion() #if NCCL_MAJOR >= 2 -ncclDataType_t test_types[ncclNumTypes] = {ncclInt8, ncclUint8, ncclInt32, ncclUint32, ncclInt64, ncclUint64, ncclHalf, ncclFloat, ncclDouble, -#if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) - ncclBfloat16 -#endif -}; -const char *test_typenames[ncclNumTypes] = {"int8", "uint8", "int32", "uint32", "int64", "uint64", "half", "float", "double", -#if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) - "bfloat16" -#endif -}; + ncclDataType_t test_types[ncclNumTypes] = { + ncclInt8, ncclUint8, ncclInt32, ncclUint32, ncclInt64, ncclUint64, ncclHalf, ncclFloat, ncclDouble + #if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) + , ncclBfloat16 + #endif + }; + const char *test_typenames[ncclNumTypes] = { + "int8", "uint8", "int32", "uint32", "int64", "uint64", "half", "float", "double" + #if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) + , "bfloat16" + #endif + }; + int test_typenum = -1; -#if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) -int test_typenum = 10; + const char *test_opnames[] = {"sum", "prod", "max", "min", "avg", "mulsum"}; + ncclRedOp_t test_ops[] = {ncclSum, ncclProd, ncclMax, ncclMin + #if NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) + , ncclAvg + #endif + #if NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) + , ncclNumOps // stand in for ncclRedOpCreatePreMulSum() created on-demand + #endif + }; + int test_opnum = -1; #else -int test_typenum = 9; -#endif - -#else -ncclDataType_t test_types[ncclNumTypes] = {ncclChar, ncclInt, ncclHalf, ncclFloat, ncclDouble, ncclInt64, ncclUint64}; -const char *test_typenames[ncclNumTypes] = {"char", "int", "half", "float", "double", "int64", "uint64"}; -int test_typenum = 7; -#endif - -#if NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) -ncclRedOp_t test_ops[ncclNumOps] = {ncclSum, ncclProd, ncclMax, ncclMin, ncclAvg}; -const char *test_opnames[ncclNumOps] = {"sum", "prod", "max", "min", "avg"}; -int test_opnum = 5; -#else -ncclRedOp_t test_ops[ncclNumOps] = {ncclSum, ncclProd, ncclMax, ncclMin}; -const char *test_opnames[ncclNumOps] = {"sum", "prod", "max", "min"}; -int test_opnum = 4; + ncclDataType_t test_types[ncclNumTypes] = {ncclChar, ncclInt, ncclHalf, ncclFloat, ncclDouble, ncclInt64, ncclUint64}; + const char *test_typenames[ncclNumTypes] = {"char", "int", "half", "float", "double", "int64", "uint64"}; + int test_typenum = 7; + const char *test_opnames[] = {"sum", "prod", "max", "min"}; + ncclRedOp_t test_ops[] = {ncclSum, ncclProd, ncclMax, ncclMin}; + int test_opnum = 4; #endif thread_local int is_main_thread = 0; @@ -265,45 +265,73 @@ template<> __device__ half ncclOpMin(half a, half b) { return __half2float(a)<__half2float(b) ? a : b; } template -__device__ T ncclPostOpIdent(T x, int n) { return x; } - +__device__ T ncclPPOpIdent(T x, int arg) { return x; } template -__device__ T ncclPostOpDiv(T x, int n) { return x/n; } +__device__ T ncclPPOpMul(T x, int arg) { return x*T(arg); } +template +__device__ T ncclPPOpDiv(T x, int arg) { return x/T(arg); } template<> -__device__ half ncclPostOpDiv(half x, int n) { return __float2half(__half2float(x)/n); } +__device__ half ncclPPOpMul(half x, int arg) { + return __float2half(__half2float(x)*float(arg)); +} +template<> +__device__ half ncclPPOpDiv(half x, int n) { + return __float2half(__half2float(x)/n); +} #if defined(__CUDA_BF16_TYPES_EXIST__) template<> -__device__ __nv_bfloat16 ncclPostOpDiv<__nv_bfloat16>(__nv_bfloat16 x, int n) { return __float2bfloat16(__bfloat162float(x)/n); } +__device__ __nv_bfloat16 ncclPPOpMul(__nv_bfloat16 x, int arg) { + return __float2bfloat16(__bfloat162float(x)*float(arg)); +} +template<> +__device__ __nv_bfloat16 ncclPPOpDiv(__nv_bfloat16 x, int n) { + return __float2bfloat16(__bfloat162float(x)/n); +} #endif -template +__host__ __device__ int preMulScalar(int rank) { + return 1 + rank%2; +} + +template __global__ void InitDataReduceKernel(T* data, const size_t N, const size_t offset, const int rep, const int nranks) { for (size_t o=blockIdx.x*blockDim.x+threadIdx.x; o(o+offset, rep, 0); + val = PreOp(val, preMulScalar(0)); for (int i=1; i(o+offset, rep, i)); + T val1 = testValue(o+offset, rep, i); + val1 = PreOp(val1, preMulScalar(i)); + val = Op(val, val1); } data[o] = PostOp(val, nranks); } } -#define KERN(type, op, postop) (void*)InitDataReduceKernel, postop > -#if NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) +#define KERN(type, op, preop, postop) (void*)InitDataReduceKernel, preop, postop > +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) #define OPS(type) \ - KERN(type, ncclOpSum, ncclPostOpIdent), \ - KERN(type, ncclOpProd, ncclPostOpIdent), \ - KERN(type, ncclOpMax, ncclPostOpIdent), \ - KERN(type, ncclOpMin, ncclPostOpIdent), \ - KERN(type, ncclOpSum/*Avg*/, ncclPostOpDiv) + KERN(type, ncclOpSum, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpProd, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpMax, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpMin, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpSum/*Avg*/, ncclPPOpIdent, ncclPPOpDiv), \ + KERN(type, ncclOpSum/*PreMulSum*/, ncclPPOpMul, ncclPPOpIdent) +#elif NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) + #define OPS(type) \ + KERN(type, ncclOpSum, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpProd, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpMax, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpMin, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpSum/*Avg*/, ncclPPOpIdent, ncclPPOpDiv) #else #define OPS(type) \ - KERN(type, ncclOpSum, ncclPostOpIdent), \ - KERN(type, ncclOpProd, ncclPostOpIdent), \ - KERN(type, ncclOpMax, ncclPostOpIdent), \ - KERN(type, ncclOpMin, ncclPostOpIdent) + KERN(type, ncclOpSum, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpProd, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpMax, ncclPPOpIdent, ncclPPOpIdent), \ + KERN(type, ncclOpMin, ncclPPOpIdent, ncclPPOpIdent) #endif -static void* const redInitDataKerns[ncclNumOps*ncclNumTypes] = { +static void* const redInitDataKerns[test_opNumMax*ncclNumTypes] = { OPS(int8_t), OPS(uint8_t), OPS(int32_t), OPS(uint32_t), OPS(int64_t), OPS(uint64_t), OPS(half), OPS(float), OPS(double), #if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) OPS(__nv_bfloat16) @@ -314,7 +342,7 @@ testResult_t InitDataReduce(void* data, const size_t count, const size_t offset, dim3 grid = { 32, 1, 1 }; dim3 block = { 256, 1, 1 }; void* args[5] = { (void*)&data, (void*)&count, (void*)&offset, (void*)&rep, (void*)&nranks }; - CUDACHECK(cudaLaunchKernel(redInitDataKerns[type*ncclNumOps+op], grid, block, args, 0, cudaStreamDefault)); + CUDACHECK(cudaLaunchKernel(redInitDataKerns[type*test_opNumMax+op], grid, block, args, 0, cudaStreamDefault)); return testSuccess; } @@ -335,7 +363,7 @@ static void* const initDataKerns[ncclNumTypes] = { (void*)InitDataKernel< float>, (void*)InitDataKernel< double>, #if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) - (void*)InitDataKernel<__nv_bfloat16>, + (void*)InitDataKernel<__nv_bfloat16> #endif }; @@ -481,7 +509,7 @@ testResult_t testStreamSynchronize(int ngpus, cudaStream_t* streams, ncclComm_t* return testSuccess; } -testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, int iter) { +testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t opIndex, int root, int in_place, int iter) { size_t count = args->nbytes / wordSize(type); // Try to change offset for each iteration so that we avoid cache effects and catch race conditions in ptrExchange @@ -499,10 +527,49 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t int rank = ((args->proc*args->nThreads + args->thread)*args->nGpus + i); char* recvBuff = ((char*)args->recvbuffs[i]) + shift; char* sendBuff = ((char*)args->sendbuffs[i]) + shift; + ncclRedOp_t op; + + if(opIndex < ncclNumOps) { + op = opIndex; + } + #if NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) + else { + union { + int8_t i8; uint8_t u8; int32_t i32; uint32_t u32; int64_t i64; uint64_t u64; + half f16; float f32; double f64; + #if defined(__CUDA_BF16_TYPES_EXIST__) + __nv_bfloat16 bf16; + #endif + }; + int scalar = preMulScalar(rank); + switch(type) { + case ncclInt8: i8 = int8_t(scalar); break; + case ncclUint8: u8 = uint8_t(scalar); break; + case ncclInt32: i32 = int32_t(scalar); break; + case ncclUint32: u32 = uint32_t(scalar); break; + case ncclInt64: i64 = int32_t(scalar); break; + case ncclUint64: u64 = uint32_t(scalar); break; + case ncclFloat16: f16 = __float2half(float(scalar)); break; + case ncclFloat32: f32 = float(scalar); break; + case ncclFloat64: f64 = double(scalar); break; + #if defined(__CUDA_BF16_TYPES_EXIST__) + case ncclBfloat16: bf16 = __float2bfloat16(float(scalar)); break; + #endif + } + NCCLCHECK(ncclRedOpCreatePreMulSum(&op, &u64, type, ncclScalarHostImmediate, args->comms[i])); + } + #endif + TESTCHECK(args->collTest->runColl( (void*)(in_place ? recvBuff + args->sendInplaceOffset*rank : sendBuff), (void*)(in_place ? recvBuff + args->recvInplaceOffset*rank : recvBuff), count, type, op, root, args->comms[i], args->streams[i])); + + #if NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) + if(opIndex >= ncclNumOps) { + NCCLCHECK(ncclRedOpDestroy(op, args->comms[i])); + } + #endif } if (args->nGpus > 1) NCCLCHECK(ncclGroupEnd()); @@ -540,7 +607,10 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t if (cudaGraphLaunches >= 1) { // Begin cuda graph capture for (int i=0; inGpus; i++) { - CUDACHECK(cudaStreamBeginCapture(args->streams[i], args->nThreads > 1 ? cudaStreamCaptureModeThreadLocal : cudaStreamCaptureModeGlobal)); + // Thread local mode is needed for: + // - Multi-thread mode + // - P2P pre-connect + CUDACHECK(cudaStreamBeginCapture(args->streams[i], cudaStreamCaptureModeThreadLocal)); } } #endif @@ -610,7 +680,7 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t if (cudaGraphLaunches >= 1) { // Begin cuda graph capture for data check for (int i=0; inGpus; i++) { - CUDACHECK(cudaStreamBeginCapture(args->streams[i], cudaStreamCaptureModeThreadLocal)); + CUDACHECK(cudaStreamBeginCapture(args->streams[i], args->nThreads > 1 ? cudaStreamCaptureModeThreadLocal : cudaStreamCaptureModeGlobal)); } } #endif @@ -777,10 +847,19 @@ int main(int argc, char* argv[]) { test_ncclVersion = NCCL_VERSION_CODE; #endif //printf("# NCCL_VERSION_CODE=%d ncclGetVersion=%d\n", NCCL_VERSION_CODE, test_ncclVersion); - if (NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) && test_ncclVersion < NCCL_VERSION(2,10,0)) { - test_opnum -= 1; // exclude ncclAvg - test_typenum -= 1; // exclude bfloat16 - } + #if NCCL_VERSION_CODE >= NCCL_VERSION(2,0,0) + test_opnum = 4; + test_typenum = 9; + if (NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) && test_ncclVersion >= NCCL_VERSION(2,10,0)) { + test_opnum++; // ncclAvg + #if defined(__CUDA_BF16_TYPES_EXIST__) + test_typenum++; // bfloat16 + #endif + } + if (NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) && test_ncclVersion >= NCCL_VERSION(2,11,0)) { + test_opnum++; // PreMulSum + } + #endif // Parse args double parsed; @@ -803,7 +882,8 @@ int main(int argc, char* argv[]) { {"blocking", required_argument, 0, 'z'}, {"cudagraph", required_argument, 0, 'G'}, {"average", required_argument, 0, 'a'}, - {"help", no_argument, 0, 'h'} + {"help", no_argument, 0, 'h'}, + {} }; while(1) { @@ -898,7 +978,9 @@ int main(int argc, char* argv[]) { "[-w,--warmup_iters ] \n\t" "[-p,--parallel_init <0/1>] \n\t" "[-c,--check <0/1>] \n\t" -#if NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) + "[-o,--op ] \n\t" +#elif NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) "[-o,--op ] \n\t" #else "[-o,--op ] \n\t" @@ -993,6 +1075,7 @@ testResult_t run() { } #ifdef MPI_SUPPORT MPI_Bcast(&ncclId, sizeof(ncclId), MPI_BYTE, 0, MPI_COMM_WORLD); + MPI_Barrier(MPI_COMM_WORLD); #endif cudaStream_t streams[nGpus*nThreads]; void* sendbuffs[nGpus*nThreads]; diff --git a/src/common.h b/src/common.h index f789c787cd..e13816f6f8 100644 --- a/src/common.h +++ b/src/common.h @@ -237,12 +237,13 @@ static size_t wordSize(ncclDataType_t type) { } extern int test_ncclVersion; // init'd with ncclGetVersion() -extern ncclDataType_t test_types[ncclNumTypes]; -extern const char *test_typenames[ncclNumTypes]; -extern ncclRedOp_t test_ops[ncclNumOps]; -extern const char *test_opnames[ncclNumOps]; +constexpr int test_opNumMax = (int)ncclNumOps + (NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) ? 1 : 0); extern int test_opnum; extern int test_typenum; +extern ncclDataType_t test_types[ncclNumTypes]; +extern const char *test_typenames[ncclNumTypes]; +extern ncclRedOp_t test_ops[]; +extern const char *test_opnames[]; static int ncclstringtotype(char *str) { for (int t=0; tsendBytes / wordSize(type); + size_t recvcount = args->expectedBytes / wordSize(type); + int nranks = args->nProcs*args->nThreads*args->nGpus; + + for (int i=0; inGpus; i++) { + int gpuid = args->localRank*args->nThreads*args->nGpus + args->thread*args->nGpus + i; + CUDACHECK(cudaSetDevice(gpuid)); + int rank = ((args->proc*args->nThreads + args->thread)*args->nGpus + i); + CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes)); + void* data = in_place ? ((char*)args->recvbuffs[i])+rank*args->sendBytes : args->sendbuffs[i]; + TESTCHECK(InitData(data, sendcount, type, rep, rank)); + CUDACHECK(cudaMemcpy(args->expected[i], args->recvbuffs[i], args->expectedBytes, cudaMemcpyDefault)); + if (rank == root) { + for (int j=0; jexpected[i])+args->sendBytes*j, sendcount, type, rep, j)); + } + } + CUDACHECK(cudaDeviceSynchronize()); + } + return testSuccess; +} + +void GatherGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw, int nranks) { + double baseBw = (double)(count * nranks * typesize) / 1.0E9 / sec; + + *algBw = baseBw; + double factor = ((double)(nranks-1))/((double)(nranks)); + *busBw = baseBw * factor; +} + +testResult_t GatherRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream) { + int nRanks; + NCCLCHECK(ncclCommCount(comm, &nRanks)); + int rank; + NCCLCHECK(ncclCommUserRank(comm, &rank)); + size_t rankOffset = count * wordSize(type); + if (count == 0) return testSuccess; + + NCCLCHECK(ncclGroupStart()); + NCCLCHECK(ncclSend(sendbuff, count, type, root, comm, stream)); + if (rank == root) { + for (int r=0; rcollTest = &gatherTest; + ncclDataType_t *run_types; + const char **run_typenames; + int type_count; + int begin_root, end_root; + + if ((int)type != -1) { + type_count = 1; + run_types = &type; + run_typenames = &typeName; + } else { + type_count = test_typenum; + run_types = test_types; + run_typenames = test_typenames; + } + + if (root != -1) { + begin_root = end_root = root; + } else { + begin_root = 0; + end_root = args->nProcs*args->nThreads*args->nGpus-1; + } + + for (int i=0; i