From c62bde3272fb94e47c94b21e53187fd77e664b49 Mon Sep 17 00:00:00 2001 From: David Addison Date: Wed, 20 Jan 2021 17:08:40 -0800 Subject: [PATCH 01/13] Do not allocate memory for expected buffer if checking disabled This allows the tests to be run with larger buffers [ROCm/rccl-tests commit: 7677f3f6081cfafb89947b675cde565c63a3f98c] --- projects/rccl-tests/src/common.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 19129d66ec..ff4e1fd857 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -559,7 +559,7 @@ testResult_t threadLaunch(struct testThread* thread) { testResult_t AllocateBuffs(void **sendbuff, size_t sendBytes, void **recvbuff, size_t recvBytes, void **expected, size_t nbytes, int nranks) { CUDACHECK(cudaMalloc(sendbuff, nbytes)); CUDACHECK(cudaMalloc(recvbuff, nbytes)); - CUDACHECK(cudaMalloc(expected, recvBytes)); + if (datacheck) CUDACHECK(cudaMalloc(expected, recvBytes)); return testSuccess; } @@ -764,7 +764,7 @@ testResult_t run() { for (int i=0; i Date: Mon, 15 Mar 2021 14:44:06 -0700 Subject: [PATCH 02/13] Add support for new datatype: bfloat16 [ROCm/rccl-tests commit: e37545e4911c210558baba789941ea7bf59db00d] --- projects/rccl-tests/src/common.cu | 41 ++++++++++++++++++++++++++++--- projects/rccl-tests/src/common.h | 3 +++ 2 files changed, 40 insertions(+), 4 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index ff4e1fd857..4589593b07 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -12,8 +12,16 @@ #include "cuda.h" #if NCCL_MAJOR >= 2 -ncclDataType_t test_types[ncclNumTypes] = {ncclInt8, ncclUint8, ncclInt32, ncclUint32, ncclInt64, ncclUint64, ncclHalf, ncclFloat, ncclDouble}; -const char *test_typenames[ncclNumTypes] = {"int8", "uint8", "int32", "uint32", "int64", "uint64", "half", "float", "double"}; +ncclDataType_t test_types[ncclNumTypes] = {ncclInt8, ncclUint8, ncclInt32, ncclUint32, ncclInt64, ncclUint64, ncclHalf, ncclFloat, ncclDouble, +#if defined(__CUDA_BF16_TYPES_EXIST__) + ncclBfloat16 +#endif +}; +const char *test_typenames[ncclNumTypes] = {"int8", "uint8", "int32", "uint32", "int64", "uint64", "half", "float", "double", +#if defined(__CUDA_BF16_TYPES_EXIST__) + "bfloat16" +#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"}; @@ -61,6 +69,9 @@ double parsesize(char *value) { double DeltaMaxValue(ncclDataType_t type) { switch(type) { case ncclHalf: return 1e-2; +#if defined(__CUDA_BF16_TYPES_EXIST__) + case ncclBfloat16: return 1e-2; +#endif case ncclFloat: return 1e-5; case ncclDouble: return 1e-12; case ncclInt: @@ -95,6 +106,12 @@ template<> __device__ float toFloat(half a) { return __half2float(a); } +#if defined(__CUDA_BF16_TYPES_EXIST__) +template<> __device__ +float toFloat(__nv_bfloat16 a) { + return __bfloat162float(a); +} +#endif template __global__ void deltaKern(void* A_, void* B_, size_t count, double* max) { @@ -128,6 +145,10 @@ void deltaKern(void* A_, void* B_, size_t count, double* max) { testResult_t CheckDelta(void* expected, void* results, size_t count, ncclDataType_t type, double* devmax) { switch (type) { +#if defined(__CUDA_BF16_TYPES_EXIST__) + case ncclBfloat16: + deltaKern<__nv_bfloat16, 512><<<1, 512>>>(results, expected, count, devmax); break; +#endif case ncclHalf: deltaKern<<<1, 512>>>(results, expected, count, devmax); break; case ncclFloat: @@ -174,6 +195,12 @@ template<> __device__ half testValue(const size_t offset, const int rep, const int rank) { return __float2half(testValue(offset, rep, rank)); } +#if defined(__CUDA_BF16_TYPES_EXIST__) +template<> +__device__ __nv_bfloat16 testValue<__nv_bfloat16>(const size_t offset, const int rep, const int rank) { + return __float2bfloat16(testValue(offset, rep, rank)); +} +#endif // Operations template @@ -210,7 +237,10 @@ __global__ void InitDataReduceKernel(T* data, const size_t N, const size_t offse #define OPS(type) KERN(type, ncclOpSum), KERN(type, ncclOpProd), KERN(type, ncclOpMax), KERN(type, ncclOpMin) static void* const redInitDataKerns[ncclNumOps*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) + 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__) + OPS(__nv_bfloat16) +#endif }; testResult_t InitDataReduce(void* data, const size_t count, const size_t offset, ncclDataType_t type, ncclRedOp_t op, const int rep, const int nranks) { @@ -236,7 +266,10 @@ static void* const initDataKerns[ncclNumTypes] = { (void*)InitDataKernel, (void*)InitDataKernel< half>, (void*)InitDataKernel< float>, - (void*)InitDataKernel< double> + (void*)InitDataKernel< double>, +#if defined(__CUDA_BF16_TYPES_EXIST__) + (void*)InitDataKernel<__nv_bfloat16>, +#endif }; template diff --git a/projects/rccl-tests/src/common.h b/projects/rccl-tests/src/common.h index 865ee258cd..c869254669 100644 --- a/projects/rccl-tests/src/common.h +++ b/projects/rccl-tests/src/common.h @@ -213,6 +213,9 @@ static size_t wordSize(ncclDataType_t type) { #endif return 1; case ncclHalf: +#if defined(__CUDA_BF16_TYPES_EXIST__) + case ncclBfloat16: +#endif //case ncclFloat16: return 2; case ncclInt: From 05f0ab10e6849e48312875383dc6c674ce4c93d5 Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Thu, 27 May 2021 09:12:52 -0700 Subject: [PATCH 03/13] Update PERFORMANCE.md [ROCm/rccl-tests commit: e12c35d84b026acc5fb573f7ac5a732430eedd32] --- projects/rccl-tests/doc/PERFORMANCE.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/rccl-tests/doc/PERFORMANCE.md b/projects/rccl-tests/doc/PERFORMANCE.md index 7cc6ecee66..21fef609af 100644 --- a/projects/rccl-tests/doc/PERFORMANCE.md +++ b/projects/rccl-tests/doc/PERFORMANCE.md @@ -46,7 +46,7 @@ A tree would do it hierarchically : `(((((i_{n-1} + i_{n-2}) + (i_{n-3} + i_{n-4})) + ... + (i_1 + i_0))))) -> o_0 -> (o_{n/2} -> (o_{3n/4} ...))` -In all cases, we need n-1 additions and n assignations for each element. Since every step is on a different rank except potentially one (the last input and the first output), +In all cases, we need n-1 additions and n assignments for each element. Since every step is on a different rank except potentially one (the last input and the first output), we need 2(n-1) data transfers (x number of elements) to perform an allReduce operation. Considering that each rank has a bandwidth to the outside world of _B_, the time to perform an allReduce operation of _S_ elements is at best : @@ -82,7 +82,7 @@ Note that here, S is the size in bytes of the total array, which for NCCL is equ ### AllGather -The AllGather operation requires only to perform the assignation part of the allReduce operation : +The AllGather operation requires only to perform the assignment part of the allReduce operation : `o_0 = o_1 = o_2 = ... = o_{n-1} = i_K` From 45c28c6c3666767af788cf96363aee54872f7952 Mon Sep 17 00:00:00 2001 From: Greg Inozemtsev Date: Wed, 2 Jun 2021 17:52:11 -0700 Subject: [PATCH 04/13] Cleanup argument error handling and messages Add error checking for minbytes and maxbytes arguments Also accept lowercase literals when parsing size arguments and print errors and usage on stderr. [ROCm/rccl-tests commit: c4de829d9131d83e4b0ca5c08cd9a8eca2dfc289] --- projects/rccl-tests/src/common.cu | 69 +++++++++++++++++++++++-------- 1 file changed, 52 insertions(+), 17 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index ff4e1fd857..25fc7dac7a 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -40,22 +40,40 @@ static int ncclroot = 0; static int parallel_init = 0; static int blocking_coll = 0; -double parsesize(char *value) { +static double parsesize(const char *value) { long long int units; double size; + char size_lit; - if (strchr(value, 'G') != NULL) { - units=1024*1024*1024; - } else if (strchr(value, 'M') != NULL) { - units=1024*1024; - } else if (strchr(value, 'K') != NULL) { - units=1024; - } else { - units=1; + int count = sscanf(value, "%lf %1s", &size, &size_lit); + + switch (count) { + case 2: + switch (size_lit) { + case 'G': + case 'g': + units = 1024*1024*1024; + break; + case 'M': + case 'm': + units = 1024*1024; + break; + case 'K': + case 'k': + units = 1024; + break; + default: + return -1.0; + }; + break; + case 1: + units = 1; + break; + default: + return -1.0; } - size = atof(value)*units; - return size; + return size * units; } double DeltaMaxValue(ncclDataType_t type) { @@ -570,6 +588,7 @@ int main(int argc, char* argv[]) { setlinebuf(stdout); // Parse args + double parsed; int longindex; static struct option longopts[] = { {"nthreads", required_argument, 0, 't'}, @@ -605,10 +624,20 @@ int main(int argc, char* argv[]) { nGpus = strtol(optarg, NULL, 0); break; case 'b': - minBytes = (size_t)parsesize(optarg); + parsed = parsesize(optarg); + if (parsed < 0) { + fprintf(stderr, "invalid size specified for 'minbytes'\n"); + return -1; + } + minBytes = (size_t)parsed; break; case 'e': - maxBytes = (size_t)parsesize(optarg); + parsed = parsesize(optarg); + if (parsed < 0) { + fprintf(stderr, "invalid size specified for 'maxbytes'\n"); + return -1; + } + maxBytes = (size_t)parsed; break; case 'i': stepBytes = strtol(optarg, NULL, 0); @@ -623,7 +652,7 @@ int main(int argc, char* argv[]) { #if NCCL_MAJOR >= 2 && NCCL_MINOR >= 2 agg_iters = (int)strtol(optarg, NULL, 0); #else - printf("Option -m not supported before NCCL 2.2. Ignoring\n"); + fprintf(stderr, "Option -m not supported before NCCL 2.2. Ignoring\n"); #endif break; case 'w': @@ -648,7 +677,7 @@ int main(int argc, char* argv[]) { blocking_coll = strtol(optarg, NULL, 0); break; case 'h': - printf("USAGE: %s \n\t" + fprintf(stderr, "USAGE: %s \n\t" "[-t,--nthreads ] \n\t" "[-g,--ngpus ] \n\t" "[-b,--minbytes ] \n\t" @@ -668,8 +697,8 @@ int main(int argc, char* argv[]) { basename(argv[0])); return 0; default: - printf("invalid option \n"); - printf("USAGE: %s \n\t" + fprintf(stderr, "invalid option \n"); + fprintf(stderr, "USAGE: %s \n\t" "[-t,--nthreads ] \n\t" "[-g,--ngpus ] \n\t" "[-b,--minbytes ] \n\t" @@ -690,6 +719,12 @@ int main(int argc, char* argv[]) { return 0; } } + if (minBytes > maxBytes) { + fprintf(stderr, "invalid sizes for 'minbytes' and 'maxbytes': %llu > %llu\n", + (unsigned long long)minBytes, + (unsigned long long)maxBytes); + return -1; + } #ifdef MPI_SUPPORT MPI_Init(&argc, &argv); #endif From a41268e26eeef01c46d23bce8d52973f37118c0b Mon Sep 17 00:00:00 2001 From: David Addison Date: Thu, 17 Jun 2021 14:08:43 -0700 Subject: [PATCH 05/13] Add support for ncclAvg operation [ROCm/rccl-tests commit: cde7e769c1879a77daddebe9da164513e030105b] --- projects/rccl-tests/src/all_gather.cu | 2 +- projects/rccl-tests/src/all_reduce.cu | 4 +- projects/rccl-tests/src/alltoall.cu | 2 +- projects/rccl-tests/src/broadcast.cu | 2 +- projects/rccl-tests/src/common.cu | 92 +++++++++++++++++++---- projects/rccl-tests/src/common.h | 5 +- projects/rccl-tests/src/reduce.cu | 4 +- projects/rccl-tests/src/reduce_scatter.cu | 4 +- 8 files changed, 90 insertions(+), 25 deletions(-) diff --git a/projects/rccl-tests/src/all_gather.cu b/projects/rccl-tests/src/all_gather.cu index f5bc44c57d..ee1d0ea0b9 100644 --- a/projects/rccl-tests/src/all_gather.cu +++ b/projects/rccl-tests/src/all_gather.cu @@ -84,7 +84,7 @@ testResult_t AllGatherRunTest(struct threadArgs* args, int root, ncclDataType_t run_types = &type; run_typenames = &typeName; } else { - type_count = ncclNumTypes; + type_count = test_typenum; run_types = test_types; run_typenames = test_typenames; } diff --git a/projects/rccl-tests/src/all_reduce.cu b/projects/rccl-tests/src/all_reduce.cu index bd8daaf0a2..52dce8993c 100644 --- a/projects/rccl-tests/src/all_reduce.cu +++ b/projects/rccl-tests/src/all_reduce.cu @@ -83,7 +83,7 @@ testResult_t AllReduceRunTest(struct threadArgs* args, int root, ncclDataType_t run_types = &type; run_typenames = &typeName; } else { - type_count = ncclNumTypes; + type_count = test_typenum; run_types = test_types; run_typenames = test_typenames; } @@ -93,7 +93,7 @@ testResult_t AllReduceRunTest(struct threadArgs* args, int root, ncclDataType_t run_ops = &op; run_opnames = &opName; } else { - op_count = ncclNumOps; + op_count = test_opnum; run_ops = test_ops; run_opnames = test_opnames; } diff --git a/projects/rccl-tests/src/alltoall.cu b/projects/rccl-tests/src/alltoall.cu index 31cfca090d..4afd3eb947 100644 --- a/projects/rccl-tests/src/alltoall.cu +++ b/projects/rccl-tests/src/alltoall.cu @@ -102,7 +102,7 @@ testResult_t AlltoAllRunTest(struct threadArgs* args, int root, ncclDataType_t t run_types = &type; run_typenames = &typeName; } else { - type_count = ncclNumTypes; + type_count = test_typenum; run_types = test_types; run_typenames = test_typenames; } diff --git a/projects/rccl-tests/src/broadcast.cu b/projects/rccl-tests/src/broadcast.cu index c62a99ff62..f7c0094864 100644 --- a/projects/rccl-tests/src/broadcast.cu +++ b/projects/rccl-tests/src/broadcast.cu @@ -92,7 +92,7 @@ testResult_t BroadcastRunTest(struct threadArgs* args, int root, ncclDataType_t run_types = &type; run_typenames = &typeName; } else { - type_count = ncclNumTypes; + type_count = test_typenum; run_types = test_types; run_typenames = test_typenames; } diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 4589593b07..1313079e79 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -11,23 +11,41 @@ #include #include "cuda.h" +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__) +#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__) +#if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) "bfloat16" #endif }; + +#if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) +int test_typenum = 10; +#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; +#endif thread_local int is_main_thread = 0; @@ -126,7 +144,7 @@ void deltaKern(void* A_, void* B_, size_t count, double* max) { if( delta > locmax ) { locmax = delta; #ifdef DEBUG_PRINT - if (delta > .1) printf("Error at %d/%ld : %f != %f\n", i, count, toFloat(A[i]), toFloat(B[i])); + if (delta > .1) printf("Error at %ld/%ld(%p) : %f != %f\n", i, count, B+i, toFloat(A[i]), toFloat(B[i])); #endif } } @@ -222,23 +240,48 @@ __device__ half ncclOpMax(half a, half b) { return __half2float(a)>__half2float( template<> __device__ half ncclOpMin(half a, half b) { return __half2float(a)<__half2float(b) ? a : b; } -template +template +__device__ T ncclPostOpIdent(T x, int n) { return x; } + +template +__device__ T ncclPostOpDiv(T x, int n) { return x/n; } +template<> +__device__ half ncclPostOpDiv(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); } +#endif + +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); for (int i=1; i(o+offset, rep, i)); } - data[o] = val; + data[o] = PostOp(val, nranks); } } -#define KERN(type, op) (void*)InitDataReduceKernel> -#define OPS(type) KERN(type, ncclOpSum), KERN(type, ncclOpProd), KERN(type, ncclOpMax), KERN(type, ncclOpMin) +#define KERN(type, op, postop) (void*)InitDataReduceKernel, postop > +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,10,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) +#else + #define OPS(type) \ + KERN(type, ncclOpSum, ncclPostOpIdent), \ + KERN(type, ncclOpProd, ncclPostOpIdent), \ + KERN(type, ncclOpMax, ncclPostOpIdent), \ + KERN(type, ncclOpMin, ncclPostOpIdent) +#endif static void* const redInitDataKerns[ncclNumOps*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__) +#if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) OPS(__nv_bfloat16) #endif }; @@ -267,7 +310,7 @@ static void* const initDataKerns[ncclNumTypes] = { (void*)InitDataKernel< half>, (void*)InitDataKernel< float>, (void*)InitDataKernel< double>, -#if defined(__CUDA_BF16_TYPES_EXIST__) +#if defined(__CUDA_BF16_TYPES_EXIST__) && NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) (void*)InitDataKernel<__nv_bfloat16>, #endif }; @@ -367,7 +410,7 @@ testResult_t testStreamSynchronize(int ngpus, cudaStream_t* streams, ncclComm_t* if (cudaErr != cudaErrorNotReady) CUDACHECK(cudaErr); #if NCCL_VERSION_CODE >= NCCL_VERSION(2,4,0) - if (comms) { + if (test_ncclVersion >= NCCL_VERSION(2,4,0) && comms) { ncclResult_t ncclAsyncErr; NCCLCHECK(ncclCommGetAsyncError(comms[i], &ncclAsyncErr)); if (ncclAsyncErr != ncclSuccess) { @@ -602,6 +645,17 @@ int main(int argc, char* argv[]) { // Make sure everyline is flushed so that we see the progress of the test setlinebuf(stdout); + #if NCCL_VERSION_CODE >= NCCL_VERSION(2,4,0) + ncclGetVersion(&test_ncclVersion); + #else + 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 + } + // Parse args int longindex; static struct option longopts[] = { @@ -653,7 +707,7 @@ int main(int argc, char* argv[]) { iters = (int)strtol(optarg, NULL, 0); break; case 'm': -#if NCCL_MAJOR >= 2 && NCCL_MINOR >= 2 +#if NCCL_MAJOR > 2 || (NCCL_MAJOR >= 2 && NCCL_MINOR >= 2) agg_iters = (int)strtol(optarg, NULL, 0); #else printf("Option -m not supported before NCCL 2.2. Ignoring\n"); @@ -693,7 +747,11 @@ 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) + "[-o,--op ] \n\t" +#else "[-o,--op ] \n\t" +#endif "[-d,--datatype ] \n\t" "[-r,--root ] \n\t" "[-z,--blocking <0/1>] \n\t" @@ -701,8 +759,8 @@ int main(int argc, char* argv[]) { basename(argv[0])); return 0; default: - printf("invalid option \n"); - printf("USAGE: %s \n\t" + if (c != 'h') printf("invalid option '%c'\n", c); + printf("USAGE: %s \n\t" "[-t,--nthreads ] \n\t" "[-g,--ngpus ] \n\t" "[-b,--minbytes ] \n\t" @@ -714,7 +772,11 @@ 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) + "[-o,--op ] \n\t" +#else "[-o,--op ] \n\t" +#endif "[-d,--datatype ] \n\t" "[-r,--root ] \n\t" "[-z,--blocking <0/1>] \n\t" @@ -899,8 +961,8 @@ testResult_t run() { // Free off CUDA allocated memory for (int i=0; i Date: Mon, 28 Jun 2021 10:12:34 -0700 Subject: [PATCH 06/13] Fixed formatting for bfloat16 support [ROCm/rccl-tests commit: 526eacadf79debe32ae0078cf2cb0fc9ed9c1baf] --- projects/rccl-tests/src/all_gather.cu | 8 ++++---- projects/rccl-tests/src/all_reduce.cu | 8 ++++---- projects/rccl-tests/src/alltoall.cu | 8 ++++---- projects/rccl-tests/src/broadcast.cu | 8 ++++---- projects/rccl-tests/src/reduce.cu | 8 ++++---- projects/rccl-tests/src/reduce_scatter.cu | 8 ++++---- 6 files changed, 24 insertions(+), 24 deletions(-) diff --git a/projects/rccl-tests/src/all_gather.cu b/projects/rccl-tests/src/all_gather.cu index ee1d0ea0b9..0b9e0cc939 100644 --- a/projects/rccl-tests/src/all_gather.cu +++ b/projects/rccl-tests/src/all_gather.cu @@ -8,15 +8,15 @@ #include "common.h" void print_header() { - PRINT("# %10s %12s %6s out-of-place in-place \n", "", "", ""); - PRINT("# %10s %12s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", + PRINT("# %10s %12s %8s out-of-place in-place \n", "", "", ""); + PRINT("# %10s %12s %8s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); - PRINT("# %10s %12s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", + PRINT("# %10s %12s %8s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", ""); } void print_line_header (size_t size, size_t count, const char *typeName, const char *opName, int root) { - PRINT("%12li %12li %6s", size, count, typeName); + PRINT("%12li %12li %8s", size, count, typeName); } void AllGatherGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { diff --git a/projects/rccl-tests/src/all_reduce.cu b/projects/rccl-tests/src/all_reduce.cu index 52dce8993c..9b6b7f02b9 100644 --- a/projects/rccl-tests/src/all_reduce.cu +++ b/projects/rccl-tests/src/all_reduce.cu @@ -8,15 +8,15 @@ #include "common.h" void print_header() { - PRINT("# %10s %12s %6s %6s out-of-place in-place \n", "", "", "", ""); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", + PRINT("# %10s %12s %8s %6s out-of-place in-place \n", "", "", "", ""); + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", ""); } void print_line_header (size_t size, size_t count, const char *typeName, const char *opName, int root) { - PRINT("%12li %12li %6s %6s", size, count, typeName, opName); + PRINT("%12li %12li %8s %6s", size, count, typeName, opName); } void AllReduceGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { diff --git a/projects/rccl-tests/src/alltoall.cu b/projects/rccl-tests/src/alltoall.cu index 4afd3eb947..865099743d 100644 --- a/projects/rccl-tests/src/alltoall.cu +++ b/projects/rccl-tests/src/alltoall.cu @@ -8,15 +8,15 @@ #include "common.h" void print_header() { - PRINT("# %10s %12s %6s %6s out-of-place in-place \n", "", "", "", ""); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", + PRINT("# %10s %12s %8s %6s out-of-place in-place \n", "", "", "", ""); + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", ""); } void print_line_header (size_t size, size_t count, const char *typeName, const char *opName, int root) { - PRINT("%12li %12li %6s %6s", size, count, typeName, opName); + PRINT("%12li %12li %8s %6s", size, count, typeName, opName); } void AlltoAllGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { diff --git a/projects/rccl-tests/src/broadcast.cu b/projects/rccl-tests/src/broadcast.cu index f7c0094864..e2b4421ac5 100644 --- a/projects/rccl-tests/src/broadcast.cu +++ b/projects/rccl-tests/src/broadcast.cu @@ -8,15 +8,15 @@ #include "common.h" void print_header() { - PRINT("# %10s %12s %6s %6s out-of-place in-place \n", "", "", "", ""); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "root", + PRINT("# %10s %12s %8s %6s out-of-place in-place \n", "", "", "", ""); + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "root", "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", ""); } void print_line_header (size_t size, size_t count, const char *typeName, const char *opName, int root) { - PRINT("%12li %12li %6s %6i", size, count, typeName, root); + PRINT("%12li %12li %8s %6i", size, count, typeName, root); } void BroadcastGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { diff --git a/projects/rccl-tests/src/reduce.cu b/projects/rccl-tests/src/reduce.cu index e40b501b7e..278768881d 100644 --- a/projects/rccl-tests/src/reduce.cu +++ b/projects/rccl-tests/src/reduce.cu @@ -8,15 +8,15 @@ #include "common.h" void print_header() { - PRINT("# %10s %12s %6s %6s out-of-place in-place \n", "", "", "", ""); - PRINT("# %10s %12s %6s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", "root", + PRINT("# %10s %12s %8s %6s out-of-place in-place \n", "", "", "", ""); + PRINT("# %10s %12s %8s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", "root", "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); - PRINT("# %10s %12s %6s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", "", + PRINT("# %10s %12s %8s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", "", "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", ""); } void print_line_header (size_t size, size_t count, const char *typeName, const char *opName, int root) { - PRINT("%12li %12li %6s %6s %6i", size, count, typeName, opName, root); + PRINT("%12li %12li %8s %6s %6i", size, count, typeName, opName, root); } void ReduceGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { diff --git a/projects/rccl-tests/src/reduce_scatter.cu b/projects/rccl-tests/src/reduce_scatter.cu index c6de434ebe..b0c4fab52e 100644 --- a/projects/rccl-tests/src/reduce_scatter.cu +++ b/projects/rccl-tests/src/reduce_scatter.cu @@ -8,15 +8,15 @@ #include "common.h" void print_header() { - PRINT("# %10s %12s %6s %6s out-of-place in-place \n", "", "", "", ""); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", + PRINT("# %10s %12s %8s %6s out-of-place in-place \n", "", "", "", ""); + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", "redop", "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); - PRINT("# %10s %12s %6s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", + PRINT("# %10s %12s %8s %6s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", "", "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", ""); } void print_line_header (size_t size, size_t count, const char *typeName, const char *opName, int root) { - PRINT("%12li %12li %6s %6s", size, count, typeName, opName); + PRINT("%12li %12li %8s %6s", size, count, typeName, opName); } void ReduceScatterGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { From e73e5a239b8f207e805160f0439b9d156fde55e2 Mon Sep 17 00:00:00 2001 From: David Addison Date: Mon, 28 Jun 2021 14:19:45 -0700 Subject: [PATCH 07/13] Added support for CUDA graph capture/replay (-G) [ROCm/rccl-tests commit: e55ad3796d710adcf72778dca02559dc6c9706bb] --- projects/rccl-tests/src/common.cu | 82 ++++++++++++++++++++++++++++++- 1 file changed, 81 insertions(+), 1 deletion(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 1313079e79..c180294644 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -65,6 +65,7 @@ static int nccltype = ncclFloat; static int ncclroot = 0; static int parallel_init = 0; static int blocking_coll = 0; +static int cudaGraphLaunches = 0; double parsesize(char *value) { long long int units; @@ -481,6 +482,15 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t Barrier(args); + cudaGraph_t graphs[args->nGpus]; + cudaGraphExec_t graphExec[args->nGpus]; + if (cudaGraphLaunches >= 1) { + // Begin cuda graph capture + for (int i=0; inGpus; i++) { + CUDACHECK(cudaStreamBeginCapture(args->streams[i], args->nThreads > 1 ? cudaStreamCaptureModeThreadLocal : cudaStreamCaptureModeGlobal)); + } + } + // Performance Benchmark auto start = std::chrono::high_resolution_clock::now(); for (int iter = 0; iter < iters; iter++) { @@ -490,11 +500,40 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t } if (agg_iters>1) NCCLCHECK(ncclGroupEnd()); } + + if (cudaGraphLaunches >= 1) { + // End cuda graph capture + for (int i=0; inGpus; i++) { + CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i)); + } + // Instantiate cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0)); + } + // Resync CPU, restart timing, launch cuda graph + Barrier(args); + start = std::chrono::high_resolution_clock::now(); + for (int l=0; lnGpus; i++) { + CUDACHECK(cudaGraphLaunch(graphExec[i], args->streams[i])); + } + } + } + TESTCHECK(completeColl(args)); auto delta = std::chrono::high_resolution_clock::now() - start; double deltaSec = std::chrono::duration_cast>(delta).count(); deltaSec = deltaSec/(iters*agg_iters); + if (cudaGraphLaunches >= 1) deltaSec = deltaSec/cudaGraphLaunches; + + if (cudaGraphLaunches >= 1) { + //destroy cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphExecDestroy(graphExec[i])); + CUDACHECK(cudaGraphDestroy(graphs[i])); + } + } double algBw, busBw; args->collTest->getBw(count, wordSize(type), deltaSec, &algBw, &busBw, args->nProcs*args->nThreads*args->nGpus); @@ -508,10 +547,41 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t // Initialize sendbuffs, recvbuffs and expected TESTCHECK(args->collTest->initData(args, type, op, root, rep, in_place)); + if (cudaGraphLaunches >= 1) { + // Begin cuda graph capture for data check + for (int i=0; inGpus; i++) { + CUDACHECK(cudaStreamBeginCapture(args->streams[i], cudaStreamCaptureModeThreadLocal)); + } + } + //test validation in single itertion, should ideally be included into the multi-iteration run TESTCHECK(startColl(args, type, op, root, in_place, 0)); + + if (cudaGraphLaunches >= 1) { + // End cuda graph capture + for (int i=0; inGpus; i++) { + CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i)); + } + // Instantiate cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0)); + } + // Launch cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphLaunch(graphExec[i], args->streams[i])); + } + } + TESTCHECK(completeColl(args)); + if (cudaGraphLaunches >= 1) { + //destroy cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphExecDestroy(graphExec[i])); + CUDACHECK(cudaGraphDestroy(graphs[i])); + } + } + TESTCHECK(CheckData(args, type, op, root, in_place, &maxDelta)); //aggregate delta from all threads and procs @@ -674,12 +744,13 @@ int main(int argc, char* argv[]) { {"datatype", required_argument, 0, 'd'}, {"root", required_argument, 0, 'r'}, {"blocking", required_argument, 0, 'z'}, + {"cudagraph", required_argument, 0, 'G'}, {"help", no_argument, 0, 'h'} }; while(1) { int c; - c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:p:c:o:d:r:z:h", longopts, &longindex); + c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:p:c:o:d:r:z:hG:", longopts, &longindex); if (c == -1) break; @@ -734,6 +805,13 @@ int main(int argc, char* argv[]) { case 'z': blocking_coll = strtol(optarg, NULL, 0); break; + case 'G': +#if (NCCL_MAJOR > 2 || (NCCL_MAJOR >= 2 && NCCL_MINOR >= 9)) && CUDART_VERSION >= 11030 + cudaGraphLaunches = strtol(optarg, NULL, 0); +#else + printf("Option -G (CUDA graph) not supported before NCCL 2.9 + CUDA 11.3. Ignoring\n"); +#endif + break; case 'h': printf("USAGE: %s \n\t" "[-t,--nthreads ] \n\t" @@ -755,6 +833,7 @@ int main(int argc, char* argv[]) { "[-d,--datatype ] \n\t" "[-r,--root ] \n\t" "[-z,--blocking <0/1>] \n\t" + "[-G,--cudagraph ] \n\t" "[-h,--help]\n", basename(argv[0])); return 0; @@ -780,6 +859,7 @@ int main(int argc, char* argv[]) { "[-d,--datatype ] \n\t" "[-r,--root ] \n\t" "[-z,--blocking <0/1>] \n\t" + "[-G,--cudagraph ] \n\t" "[-h,--help]\n", basename(argv[0])); return 0; From d30e35f15086054b861745b615f832a73c88447b Mon Sep 17 00:00:00 2001 From: David Addison Date: Mon, 28 Jun 2021 16:49:10 -0700 Subject: [PATCH 08/13] Added new tests: scatter, sendrecv, hypercube [ROCm/rccl-tests commit: 9dae3d3a37a7505a9eb0622be4268e2d2a3cb5f9] --- projects/rccl-tests/src/Makefile | 4 +- projects/rccl-tests/src/hypercube.cu | 124 ++++++++++++++++++++++++++ projects/rccl-tests/src/scatter.cu | 125 ++++++++++++++++++++++++++ projects/rccl-tests/src/sendrecv.cu | 127 +++++++++++++++++++++++++++ 4 files changed, 378 insertions(+), 2 deletions(-) create mode 100644 projects/rccl-tests/src/hypercube.cu create mode 100644 projects/rccl-tests/src/scatter.cu create mode 100644 projects/rccl-tests/src/sendrecv.cu diff --git a/projects/rccl-tests/src/Makefile b/projects/rccl-tests/src/Makefile index 52169bb3e1..26e653e7d6 100644 --- a/projects/rccl-tests/src/Makefile +++ b/projects/rccl-tests/src/Makefile @@ -1,5 +1,5 @@ # -# Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved. +# Copyright (c) 2015-2021, NVIDIA CORPORATION. All rights reserved. # # See LICENSE.txt for license information # @@ -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 +BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall scatter sendrecv hypercube BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf) build: ${BIN_FILES} diff --git a/projects/rccl-tests/src/hypercube.cu b/projects/rccl-tests/src/hypercube.cu new file mode 100644 index 0000000000..142f1a6359 --- /dev/null +++ b/projects/rccl-tests/src/hypercube.cu @@ -0,0 +1,124 @@ +/************************************************************************* + * Copyright (c) 2016-2021, NVIDIA CORPORATION. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +#include "cuda_runtime.h" +#include "common.h" + +#define ALIGN 4 + +void print_header() { + PRINT("# %10s %12s %8s out-of-place in-place \n", "", "", ""); + PRINT("# %10s %12s %8s %7s %6s %6s %5s %7s %6s %6s %5s\n", "size", "count", "type", + "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); + PRINT("# %10s %12s %8s %7s %6s %6s %5s %7s %6s %6s %5s\n", "(B)", "(elements)", "", + "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", ""); +} + +void print_line_header (size_t size, size_t count, const char *typeName, const char *opName, int root) { + PRINT("%12li %12li %8s", size, count, typeName); +} + +void HyperCubeGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { + size_t base = (count/(ALIGN*nranks))*ALIGN; + *sendcount = base; + *recvcount = base*nranks; + *sendInplaceOffset = base; + *recvInplaceOffset = 0; + *paramcount = base; +} + +testResult_t HyperCubeInitData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int rep, int in_place) { + size_t sendcount = args->sendBytes / 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)); + for (int j=0; jexpected[i])+args->sendBytes*j, sendcount, type, rep, j)); + } + CUDACHECK(cudaDeviceSynchronize()); + } + return testSuccess; +} + +void HyperCubeGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw, int nranks) { + double baseBw = (double)(count * typesize * (nranks - 1)) / 1.0E9 / sec; + + *algBw = baseBw; + double factor = 1; + *busBw = baseBw * factor; +} + +testResult_t HyperCubeRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream) { + char* sbuff = (char*)sendbuff; + char* rbuff = (char*)recvbuff; + int nRanks; + NCCLCHECK(ncclCommCount(comm, &nRanks)); + int rank; + NCCLCHECK(ncclCommUserRank(comm, &rank)); + size_t rankSize = count * wordSize(type); + + if (rbuff+rank*rankSize != sbuff) CUDACHECK(cudaMemcpyAsync(rbuff+rank*rankSize, sbuff, rankSize, cudaMemcpyDeviceToDevice, stream)); + + // Hypercube AllGather + for (int mask=1; maskcollTest = &hyperCubeTest; + ncclDataType_t *run_types; + const char **run_typenames; + int type_count; + + 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; + } + + for (int i=0; isendBytes / wordSize(type); + size_t recvcount = args->expectedBytes / wordSize(type); + + 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 ? args->recvbuffs[i] : args->sendbuffs[i]; + if (rank == root) TESTCHECK(InitData(data, sendcount, type, rep, rank)); + TESTCHECK(InitData(args->expected[i], recvcount, type, rep+rank*recvcount, root)); + CUDACHECK(cudaDeviceSynchronize()); + } + return testSuccess; +} + +void ScatterGetBw(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 ScatterRunColl(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()); + if (rank == root) { + for (int r=0; rcollTest = &scatterTest; + 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; isendBytes / 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 ? args->recvbuffs[i] : args->sendbuffs[i]; + TESTCHECK(InitData(data, sendcount, type, rep, rank)); + int peer = (rank-1+nranks)%nranks; + TESTCHECK(InitData(args->expected[i], recvcount, type, rep, peer)); + CUDACHECK(cudaDeviceSynchronize()); + } + // We don't support in-place sendrecv + args->reportErrors = in_place ? 0 : 1; + return testSuccess; +} + +void SendRecvGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw, int nranks) { + double baseBw = (double)(count * typesize) / 1.0E9 / sec; + + *algBw = baseBw; + double factor = 1; + *busBw = baseBw * factor; +} + +testResult_t SendRecvRunColl(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)); + int recvPeer = (rank-1+nRanks) % nRanks; + int sendPeer = (rank+1) % nRanks; + + NCCLCHECK(ncclGroupStart()); + NCCLCHECK(ncclSend(sendbuff, count, type, sendPeer, comm, stream)); + NCCLCHECK(ncclRecv(recvbuff, count, type, recvPeer, comm, stream)); + NCCLCHECK(ncclGroupEnd()); + return testSuccess; +} + +struct testColl sendRecvTest = { + "SendRecv", + SendRecvGetCollByteCount, + SendRecvInitData, + SendRecvGetBw, + SendRecvRunColl +}; + +void SendRecvGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) { + size_t paramcount, sendInplaceOffset, recvInplaceOffset; + SendRecvGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks); +} + +testResult_t SendRecvRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) { + args->collTest = &sendRecvTest; + ncclDataType_t *run_types; + ncclRedOp_t *run_ops; + const char **run_typenames, **run_opnames; + int type_count, op_count; + + 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 ((int)op != -1) { + op_count = 1; + run_ops = &op; + run_opnames = &opName; + } else { + op_count = test_opnum; + run_ops = test_ops; + run_opnames = test_opnames; + } + + for (int i=0; i Date: Mon, 28 Jun 2021 18:23:12 -0700 Subject: [PATCH 09/13] Resync with changes in gitilab-master code [ROCm/rccl-tests commit: 1ae8cdc315d10c3f65764a1915e0fb0f1563d893] --- projects/rccl-tests/src/common.cu | 81 ++++++++++++++----------------- projects/rccl-tests/src/common.h | 5 +- 2 files changed, 40 insertions(+), 46 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index c180294644..7aad2c1868 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -67,6 +67,8 @@ static int parallel_init = 0; static int blocking_coll = 0; static int cudaGraphLaunches = 0; +#define NUM_BLOCKS 32 + double parsesize(char *value) { long long int units; double size; @@ -137,9 +139,9 @@ void deltaKern(void* A_, void* B_, size_t count, double* max) { const T* A = (const T*)A_; const T* B = (const T*)B_; __shared__ double temp[BSIZE]; - int tid = threadIdx.x; + int tid = blockIdx.x*blockDim.x + threadIdx.x; double locmax = 0.0; - for(int i=tid; i locmax ) { @@ -150,6 +152,7 @@ void deltaKern(void* A_, void* B_, size_t count, double* max) { } } + tid = threadIdx.x; temp[tid] = locmax; for(int stride = BSIZE/2; stride > 1; stride>>=1) { __syncthreads(); @@ -158,38 +161,38 @@ void deltaKern(void* A_, void* B_, size_t count, double* max) { } __syncthreads(); if( threadIdx.x == 0) - *max = temp[0] > temp[1] ? temp[0] : temp[1]; + max[blockIdx.x] = temp[0] > temp[1] ? temp[0] : temp[1]; } - -testResult_t CheckDelta(void* expected, void* results, size_t count, ncclDataType_t type, double* devmax) { +testResult_t CheckDelta(void* results, void* expected, size_t count, ncclDataType_t type, double* devmax) { switch (type) { #if defined(__CUDA_BF16_TYPES_EXIST__) case ncclBfloat16: - deltaKern<__nv_bfloat16, 512><<<1, 512>>>(results, expected, count, devmax); break; + deltaKern<__nv_bfloat16, 512><<>>(results, expected, count, devmax); break; #endif case ncclHalf: - deltaKern<<<1, 512>>>(results, expected, count, devmax); break; + deltaKern<<>>(results, expected, count, devmax); break; case ncclFloat: - deltaKern<<<1, 512>>>(results, expected, count, devmax); break; + deltaKern<<>>(results, expected, count, devmax); break; case ncclDouble: - deltaKern<<<1, 512>>>(results, expected, count, devmax); break; + deltaKern<<>>(results, expected, count, devmax); break; case ncclChar: #if NCCL_MAJOR >= 2 case ncclUint8: #endif - deltaKern<<<1, 512>>>(results, expected, count, devmax); break; + deltaKern<<>>(results, expected, count, devmax); break; case ncclInt: #if NCCL_MAJOR >= 2 case ncclUint32: #endif - deltaKern<<<1, 512>>>(results, expected, count, devmax); break; + deltaKern<<>>(results, expected, count, devmax); break; case ncclInt64: case ncclUint64: - deltaKern<<<1, 512>>>(results, expected, count, devmax); break; + deltaKern<<>>(results, expected, count, devmax); break; } CUDACHECK(cudaDeviceSynchronize()); + for (int i=1; isendBytes, args->expectedBytes); - size_t shift = (totalnbytes * iter) % args->maxbytes; - if (shift + totalnbytes > args->maxbytes) shift = 0; + size_t steps = totalnbytes ? args->maxbytes / totalnbytes : 1; + size_t shift = totalnbytes * (iter % steps); if (args->nGpus > 1) NCCLCHECK(ncclGroupStart()); for (int i = 0; i < args->nGpus; i++) { @@ -475,6 +478,10 @@ testResult_t completeColl(struct threadArgs* args) { testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place) { size_t count = args->nbytes / wordSize(type); + if (datacheck) { + // Initialize sendbuffs, recvbuffs and expected + TESTCHECK(args->collTest->initData(args, type, op, root, 99, in_place)); + } // Sync TESTCHECK(startColl(args, type, op, root, in_place, 0)); @@ -598,10 +605,10 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t } double timeUsec = deltaSec*1.0E6; - char timeStr[10]; + char timeStr[100]; if (timeUsec > 10000.0) { sprintf(timeStr, "%7.0f", timeUsec); - } else if (timeUsec > 100.0) { + } else if (timeUsec >= 100.0) { sprintf(timeStr, "%7.1f", timeUsec); } else { sprintf(timeStr, "%7.2f", timeUsec); @@ -812,31 +819,6 @@ int main(int argc, char* argv[]) { printf("Option -G (CUDA graph) not supported before NCCL 2.9 + CUDA 11.3. Ignoring\n"); #endif break; - case 'h': - printf("USAGE: %s \n\t" - "[-t,--nthreads ] \n\t" - "[-g,--ngpus ] \n\t" - "[-b,--minbytes ] \n\t" - "[-e,--maxbytes ] \n\t" - "[-i,--stepbytes ] \n\t" - "[-f,--stepfactor ] \n\t" - "[-n,--iters ] \n\t" - "[-m,--agg_iters ] \n\t" - "[-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) - "[-o,--op ] \n\t" -#else - "[-o,--op ] \n\t" -#endif - "[-d,--datatype ] \n\t" - "[-r,--root ] \n\t" - "[-z,--blocking <0/1>] \n\t" - "[-G,--cudagraph ] \n\t" - "[-h,--help]\n", - basename(argv[0])); - return 0; default: if (c != 'h') printf("invalid option '%c'\n", c); printf("USAGE: %s \n\t" @@ -868,7 +850,8 @@ int main(int argc, char* argv[]) { #ifdef MPI_SUPPORT MPI_Init(&argc, &argv); #endif - return run(); + TESTCHECK(run()); + return 0; } testResult_t run() { @@ -900,6 +883,7 @@ testResult_t run() { #define MAX_LINE 2048 char line[MAX_LINE]; int len = 0; + size_t maxMem = ~0; for (int i=0; i memMaxBytes) { + maxBytes = memMaxBytes; + if (proc == 0) printf("#\n# Reducing maxBytes to %ld due to memory limitation\n", maxBytes); + } + ncclUniqueId ncclId; if (proc == 0) { NCCLCHECK(ncclGetUniqueId(&ncclId)); @@ -963,7 +956,7 @@ testResult_t run() { int errors[nThreads]; double bw[nThreads]; double* delta; - CUDACHECK(cudaHostAlloc(&delta, sizeof(double)*nThreads, cudaHostAllocPortable | cudaHostAllocMapped)); + CUDACHECK(cudaHostAlloc(&delta, sizeof(double)*nThreads*NUM_BLOCKS, cudaHostAllocPortable | cudaHostAllocMapped)); int bw_count[nThreads]; for (int t=0; t Date: Wed, 30 Jun 2021 19:36:07 -0700 Subject: [PATCH 10/13] Added new option to report average iteration time [ROCm/rccl-tests commit: 1dfc76ecccb73f7e7336730663e5901d6d1600ef] --- projects/rccl-tests/src/common.cu | 32 ++++++++++++++++++++++++++++++- 1 file changed, 31 insertions(+), 1 deletion(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 7aad2c1868..d9f036879e 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -66,6 +66,10 @@ static int ncclroot = 0; static int parallel_init = 0; static int blocking_coll = 0; static int cudaGraphLaunches = 0; +#ifdef MPI_SUPPORT +// Report average iteration time: (0=RANK0,1=AVG,2=MIN,3=MAX) +static int average = 1; +#endif #define NUM_BLOCKS 32 @@ -533,6 +537,23 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t double deltaSec = std::chrono::duration_cast>(delta).count(); deltaSec = deltaSec/(iters*agg_iters); if (cudaGraphLaunches >= 1) deltaSec = deltaSec/cudaGraphLaunches; +#ifdef MPI_SUPPORT + switch (average) { + case 1: + // Calculate the average time across all ranks + MPI_Allreduce(MPI_IN_PLACE, &deltaSec, 1, MPI_DOUBLE, MPI_SUM, MPI_COMM_WORLD); + deltaSec = deltaSec/(args->nProcs*args->nThreads*args->nGpus); + break; + case 2: + // Obtain the minimum time across all ranks + MPI_Allreduce(MPI_IN_PLACE, &deltaSec, 1, MPI_DOUBLE, MPI_MIN, MPI_COMM_WORLD); + break; + case 3: + // Obtain the maximum time across all ranks + MPI_Allreduce(MPI_IN_PLACE, &deltaSec, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD); + break; + } +#endif if (cudaGraphLaunches >= 1) { //destroy cuda graph @@ -752,12 +773,13 @@ int main(int argc, char* argv[]) { {"root", required_argument, 0, 'r'}, {"blocking", required_argument, 0, 'z'}, {"cudagraph", required_argument, 0, 'G'}, + {"average", required_argument, 0, 'a'}, {"help", no_argument, 0, 'h'} }; while(1) { int c; - c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:p:c:o:d:r:z:hG:", longopts, &longindex); + c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:p:c:o:d:r:z:hG:a:", longopts, &longindex); if (c == -1) break; @@ -819,6 +841,11 @@ int main(int argc, char* argv[]) { printf("Option -G (CUDA graph) not supported before NCCL 2.9 + CUDA 11.3. Ignoring\n"); #endif break; +#ifdef MPI_SUPPORT + case 'a': + average = (int)strtol(optarg, NULL, 0); + break; +#endif default: if (c != 'h') printf("invalid option '%c'\n", c); printf("USAGE: %s \n\t" @@ -842,6 +869,9 @@ int main(int argc, char* argv[]) { "[-r,--root ] \n\t" "[-z,--blocking <0/1>] \n\t" "[-G,--cudagraph ] \n\t" +#ifdef MPI_SUPPORT + "[-a,--average <0/1/2/3> report average iteration time <0=RANK0/1=AVG/2=MIN/3=MAX>] \n\t" +#endif "[-h,--help]\n", basename(argv[0])); return 0; From ea6eec9e8009491bfbbcfaee3f60c4aadd453598 Mon Sep 17 00:00:00 2001 From: David Addison Date: Tue, 6 Jul 2021 14:47:50 -0700 Subject: [PATCH 11/13] Updated with new command line arguments [ROCm/rccl-tests commit: 11cff17a04e268ea0a82cc8517fdcfde3414280e] --- projects/rccl-tests/README.md | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/rccl-tests/README.md b/projects/rccl-tests/README.md index 791bed2599..bff6433b89 100644 --- a/projects/rccl-tests/README.md +++ b/projects/rccl-tests/README.md @@ -52,19 +52,21 @@ All tests support the same set of arguments : * `-i,--stepbytes ` fixed increment between sizes. Default : (max-min)/10. * `-f,--stepfactor ` multiplication factor between sizes. Default : disabled. * NCCL operations arguments - * `-o,--op ` Specify which reduction operation to perform. Only relevant for reduction operations like Allreduce, Reduce or ReduceScatter. Default : Sum. + * `-o,--op ` Specify which reduction operation to perform. Only relevant for reduction operations like Allreduce, Reduce or ReduceScatter. Default : Sum. * `-d,--datatype ` Specify which datatype to use. Default : Float. * `-r,--root ` Specify which root to use. Only for operations with a root like broadcast or reduce. Default : 0. * Performance * `-n,--iters ` number of iterations. Default : 20. * `-w,--warmup_iters ` number of warmup iterations (not timed). Default : 5. * `-m,--agg_iters ` number of operations to aggregate together in each iteration. Default : 1. + * `-a,--average <0/1/2/3>` Report performance as an average across all ranks (MPI=1 only). <0=Rank0,1=Avg,2=Min,3=Max>. Default : 1. * Test operation * `-p,--parallel_init <0/1>` use threads to initialize NCCL in parallel. Default : 0. * `-c,--check <0/1>` check correctness of results. This can be quite slow on large numbers of GPUs. Default : 1. * `-z,--blocking <0/1>` Make NCCL collective blocking, i.e. have CPUs wait and sync after each collective. Default : 0. + * `-G,--cudagraph ` Capture iterations as a CUDA graph and then replay specified number of times. Default : 0. ## Copyright -NCCL tests are provided under the BSD license. All source code and accompanying documentation is copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved. +NCCL tests are provided under the BSD license. All source code and accompanying documentation is copyright (c) 2016-2021, NVIDIA CORPORATION. All rights reserved. From d3061dc2a9f9a960f1bb1b80b8964aed608733e0 Mon Sep 17 00:00:00 2001 From: David Addison Date: Thu, 8 Jul 2021 16:42:40 -0700 Subject: [PATCH 12/13] Fix issues with MPI_Allreduce and multi-threaded tests [ROCm/rccl-tests commit: 547e119d350a8ad9034c4e75c8664e62c60bf599] --- projects/rccl-tests/src/common.cu | 71 ++++++++++++++++--------------- projects/rccl-tests/src/common.h | 4 +- 2 files changed, 38 insertions(+), 37 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index d4ee519107..4768b9be3e 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -356,12 +356,9 @@ testResult_t InitData(void* data, const size_t count, ncclDataType_t type, const return testSuccess; } -void Barrier(struct threadArgs* args) -{ +void Barrier(struct threadArgs* args) { while (args->barrier[args->barrier_idx] != args->thread) pthread_yield(); - args->barrier[args->barrier_idx] = args->thread + 1; - if (args->thread+1 == args->nThreads) { #ifdef MPI_SUPPORT MPI_Barrier(MPI_COMM_WORLD); @@ -370,7 +367,35 @@ void Barrier(struct threadArgs* args) } else { while (args->barrier[args->barrier_idx]) pthread_yield(); } + args->barrier_idx=!args->barrier_idx; +} +// Inter-thread/process barrier+allreduce +void Allreduce(struct threadArgs* args, double* value, int average) { + while (args->barrier[args->barrier_idx] != args->thread) pthread_yield(); + double val = *value; + if (args->thread > 0) { + double val2 = args->reduce[args->barrier_idx]; + if (average == 1) val += val2; + if (average == 2) val = std::min(val, val2); + if (average == 3) val = std::max(val, val2); + } + if (average || args->thread == 0) args->reduce[args->barrier_idx] = val; + args->barrier[args->barrier_idx] = args->thread + 1; + if (args->thread+1 == args->nThreads) { +#ifdef MPI_SUPPORT + if (average != 0) { + MPI_Op op = average == 1 ? MPI_SUM : average == 2 ? MPI_MIN : MPI_MAX; + MPI_Allreduce(MPI_IN_PLACE, (void*)&args->reduce[args->barrier_idx], 1, MPI_DOUBLE, op, MPI_COMM_WORLD); + } +#endif + if (average == 1) args->reduce[args->barrier_idx] /= args->nProcs*args->nThreads; + args->reduce[1-args->barrier_idx] = 0; + args->barrier[args->barrier_idx] = 0; + } else { + while (args->barrier[args->barrier_idx]) pthread_yield(); + } + *value = args->reduce[args->barrier_idx]; args->barrier_idx=!args->barrier_idx; } @@ -383,7 +408,7 @@ testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t NCCLCHECK(ncclCommCuDevice(args->comms[i], &device)); CUDACHECK(cudaSetDevice(device)); void *data = in_place ? ((void *)((uintptr_t)args->recvbuffs[i] + args->recvInplaceOffset*rank)) : args->recvbuffs[i]; - TESTCHECK(CheckDelta(data , args->expected[i], count, type, args->delta)); + TESTCHECK(CheckDelta(data , args->expected[i], count, type, args->deltaHost)); maxDelta = std::max(*(args->deltaHost), maxDelta); #ifdef DEBUG_PRINT @@ -555,23 +580,7 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t double deltaSec = std::chrono::duration_cast>(delta).count(); deltaSec = deltaSec/(iters*agg_iters); if (cudaGraphLaunches >= 1) deltaSec = deltaSec/cudaGraphLaunches; -#ifdef MPI_SUPPORT - switch (average) { - case 1: - // Calculate the average time across all ranks - MPI_Allreduce(MPI_IN_PLACE, &deltaSec, 1, MPI_DOUBLE, MPI_SUM, MPI_COMM_WORLD); - deltaSec = deltaSec/(args->nProcs*args->nThreads*args->nGpus); - break; - case 2: - // Obtain the minimum time across all ranks - MPI_Allreduce(MPI_IN_PLACE, &deltaSec, 1, MPI_DOUBLE, MPI_MIN, MPI_COMM_WORLD); - break; - case 3: - // Obtain the maximum time across all ranks - MPI_Allreduce(MPI_IN_PLACE, &deltaSec, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD); - break; - } -#endif + Allreduce(args, &deltaSec, average); if (cudaGraphLaunches >= 1) { //destroy cuda graph @@ -631,21 +640,12 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t TESTCHECK(CheckData(args, type, op, root, in_place, &maxDelta)); //aggregate delta from all threads and procs - Barrier(args); - if (args->thread == 0) { - for (int i=1; inThreads; i++) { - maxDelta += args->deltaThreads[i]; - } -#ifdef MPI_SUPPORT - MPI_Allreduce(MPI_IN_PLACE, &maxDelta, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD); -#endif - } - Barrier(args); + Allreduce(args, &maxDelta, 3); } double timeUsec = deltaSec*1.0E6; char timeStr[100]; - if (timeUsec > 10000.0) { + if (timeUsec >= 10000.0) { sprintf(timeStr, "%7.0f", timeUsec); } else if (timeUsec >= 100.0) { sprintf(timeStr, "%7.1f", timeUsec); @@ -875,6 +875,7 @@ int main(int argc, char* argv[]) { average = (int)strtol(optarg, NULL, 0); break; #endif + case 'h': default: if (c != 'h') printf("invalid option '%c'\n", c); printf("USAGE: %s \n\t" @@ -1033,6 +1034,7 @@ testResult_t run() { int* sync = (int*)calloc(2, sizeof(int)); int* barrier = (int*)calloc(2, sizeof(int)); + double* reduce = (double*)calloc(2, sizeof(double)); struct testThread threads[nThreads]; memset(threads, 0, sizeof(struct testThread)*nThreads); @@ -1058,11 +1060,10 @@ testResult_t run() { threads[t].args.barrier = (volatile int*)barrier; threads[t].args.barrier_idx = 0; + threads[t].args.reduce = (volatile double*)reduce; threads[t].args.sync = (volatile int*)sync; threads[t].args.sync_idx = 0; - threads[t].args.deltaThreads = delta; threads[t].args.deltaHost = (delta + t*NUM_BLOCKS); - threads[t].args.delta = delta; threads[t].args.errors=errors+t; threads[t].args.bw=bw+t; threads[t].args.bw_count=bw_count+t; diff --git a/projects/rccl-tests/src/common.h b/projects/rccl-tests/src/common.h index 44b298dfd2..f789c787cd 100644 --- a/projects/rccl-tests/src/common.h +++ b/projects/rccl-tests/src/common.h @@ -8,6 +8,7 @@ #include "nccl.h" #include +#include #include #include #ifdef MPI_SUPPORT @@ -116,11 +117,10 @@ struct threadArgs { int sync_idx; volatile int* barrier; int barrier_idx; + volatile double* reduce; int syncRank; int syncNranks; - double* deltaThreads; double* deltaHost; - double* delta; int* errors; double* bw; int* bw_count; From 6719794fc8c346f279fb867a90ea011b62c63e84 Mon Sep 17 00:00:00 2001 From: David Addison Date: Mon, 12 Jul 2021 11:43:57 -0700 Subject: [PATCH 13/13] Removed MPI_SUPPORT conditional compilation of average flag [ROCm/rccl-tests commit: b9f90d12a906a0dc5e49f1ede6a52c7779289c01] --- projects/rccl-tests/src/common.cu | 6 ------ 1 file changed, 6 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 4768b9be3e..c343342ffa 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -66,10 +66,8 @@ static int ncclroot = 0; static int parallel_init = 0; static int blocking_coll = 0; static int cudaGraphLaunches = 0; -#ifdef MPI_SUPPORT // Report average iteration time: (0=RANK0,1=AVG,2=MIN,3=MAX) static int average = 1; -#endif #define NUM_BLOCKS 32 @@ -870,11 +868,9 @@ int main(int argc, char* argv[]) { printf("Option -G (CUDA graph) not supported before NCCL 2.9 + CUDA 11.3. Ignoring\n"); #endif break; -#ifdef MPI_SUPPORT case 'a': average = (int)strtol(optarg, NULL, 0); break; -#endif case 'h': default: if (c != 'h') printf("invalid option '%c'\n", c); @@ -899,9 +895,7 @@ int main(int argc, char* argv[]) { "[-r,--root ] \n\t" "[-z,--blocking <0/1>] \n\t" "[-G,--cudagraph ] \n\t" -#ifdef MPI_SUPPORT "[-a,--average <0/1/2/3> report average iteration time <0=RANK0/1=AVG/2=MIN/3=MAX>] \n\t" -#endif "[-h,--help]\n", basename(argv[0])); return 0;