From af08abefb06f791eb6f11d794c0080db96177649 Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Thu, 19 Mar 2020 10:18:39 -0700 Subject: [PATCH] Add gather and scatter test [ROCm/rccl-tests commit: 6932a583e74c7685fcc4e7206e367d917f3e0485] --- projects/rccl-tests/src/Makefile | 2 +- projects/rccl-tests/src/common.cu | 21 +++-- projects/rccl-tests/src/gather.cu | 127 ++++++++++++++++++++++++++++ projects/rccl-tests/src/scatter.cu | 129 +++++++++++++++++++++++++++++ 4 files changed, 269 insertions(+), 10 deletions(-) create mode 100644 projects/rccl-tests/src/gather.cu create mode 100644 projects/rccl-tests/src/scatter.cu diff --git a/projects/rccl-tests/src/Makefile b/projects/rccl-tests/src/Makefile index b02d6e886d..e109b8c3ae 100644 --- a/projects/rccl-tests/src/Makefile +++ b/projects/rccl-tests/src/Makefile @@ -54,7 +54,7 @@ HIPLDFLAGS += $(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 gather scatter BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf) build: ${BIN_FILES} diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index d14c286125..908a69b9c5 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -16,6 +16,8 @@ #include #include +//#define DEBUG_PRINT + #if NCCL_MAJOR >= 2 #if RCCL_BFLOAT16 == 1 ncclDataType_t test_types[ncclNumTypes] = {ncclInt8, ncclUint8, ncclInt32, ncclUint32, ncclInt64, ncclUint64, ncclHalf, ncclFloat, ncclDouble, ncclBfloat16}; @@ -326,6 +328,8 @@ testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t for (int i=0; inGpus; i++) { int device; int rank = ((args->proc*args->nThreads + args->thread)*args->nGpus + i); + if (rank != root && strcmp(args->collTest->name, "Gather") == 0) + continue; NCCLCHECK(ncclCommCuDevice(args->comms[i], &device)); HIPCHECK(hipSetDevice(device)); void *data = in_place ? ((void *)((uintptr_t)args->recvbuffs[i] + args->recvInplaceOffset*rank)) : args->recvbuffs[i]; @@ -333,25 +337,24 @@ testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t maxDelta = std::max(*(args->deltaHost), maxDelta); #ifdef DEBUG_PRINT - if (rank == 0) { + //if (rank == 0) { int *expectedHost = (int *)malloc(args->expectedBytes); int *dataHost = (int *)malloc(args->expectedBytes); - hipMemcpy(expectedHost, args->expected[0], args->expectedBytes, hipMemcpyDeviceToHost); - printf("\n Expected: "); + hipMemcpy(expectedHost, args->expected[rank], args->expectedBytes, hipMemcpyDeviceToHost); + printf("\n Rank [%d] Expected: ", rank); for(int j=0; jexpectedBytes/sizeof(int); j++) { printf("%d:%d ", j, expectedHost[j]); } - printf("\n"); - hipMemcpy(dataHost, data, args->expectedBytes, hipMemcpyDeviceToHost); - printf("\n Actual: "); + printf("\n Rank [%d] Actual: ", rank); for (int j=0; jexpectedBytes/sizeof(int); j++) { printf("%d:%d ", j, dataHost[j]); } printf("\n"); - free(temp); - } + free(dataHost); + free(expectedHost); + //} #endif } double nranks = args->nProcs*args->nThreads*args->nGpus; @@ -571,7 +574,7 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char* TESTCHECK(completeColl(args)); for (size_t iter = 0; iter < stress_cycles; iter++) { - if (iter > 0) PRINT("# Testing %ld cycle.\n", iter+1); + if (iter > 0) PRINT("# Testing %lu cycle.\n", iter+1); // Benchmark for (size_t size = args->minbytes; size<=args->maxbytes; size = ((args->stepfactor > 1) ? size*args->stepfactor : size+args->stepbytes)) { setupArgs(size, type, args); diff --git a/projects/rccl-tests/src/gather.cu b/projects/rccl-tests/src/gather.cu new file mode 100644 index 0000000000..5230fdc81d --- /dev/null +++ b/projects/rccl-tests/src/gather.cu @@ -0,0 +1,127 @@ +/************************************************************************* + * Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +#include +#include "common.h" + +//#define DEBUG_PRINT + +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", + "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); + PRINT("# %10s %12s %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", size, count, typeName); +} + +void GatherGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { + *sendcount = count/nranks; + *recvcount = (count/nranks)*nranks; + *sendInplaceOffset = count/nranks; + *recvInplaceOffset = 0; + *paramcount = *sendcount; +} + +testResult_t GatherInitData(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; + HIPCHECK(hipSetDevice(gpuid)); + int rank = ((args->proc*args->nThreads + args->thread)*args->nGpus + i); + HIPCHECK(hipMemset(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)); +#ifdef DEBUG_PRINT + int *dataHost = (int *)malloc(args->sendBytes); + hipMemcpy(dataHost, data, args->sendBytes, hipMemcpyDeviceToHost); + printf("\n Rank [%d] Init: ", rank); + for (int j=0; jsendBytes/sizeof(int); j++) { + printf("%d:%d ", j, dataHost[j]); + } + printf("\n"); + free(dataHost); +#endif + for (int j=0; jexpected[i])+args->sendBytes*j, sendcount, type, rep, j)); + } + HIPCHECK(hipDeviceSynchronize()); + } + return testSuccess; +} + +void GatherGetBw(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 GatherRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream) { + int nRanks; + NCCLCHECK(ncclCommCount(comm, &nRanks)); + size_t rankOffset = count * wordSize(type); + if (count == 0) return testSuccess; + + int rank; + NCCLCHECK(ncclCommUserRank(comm, &rank)); + NCCLCHECK(ncclGroupStart()); + if (rank == root) { + for (int r=0; rcollTest = &gatherTest; + 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 = ncclNumTypes; + run_types = test_types; + run_typenames = test_typenames; + } + + for (int i=0; i +#include "common.h" + +//#define DEBUG_PRINT + +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", + "time", "algbw", "busbw", "error", "time", "algbw", "busbw", "error"); + PRINT("# %10s %12s %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", size, count, typeName, opName); +} + +void ScatterGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { + *sendcount = (count/nranks)*nranks; + *recvcount = count/nranks; + *sendInplaceOffset = 0; + *recvInplaceOffset = 0; + *paramcount = *recvcount; +} + +testResult_t ScatterInitData(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; + HIPCHECK(hipSetDevice(gpuid)); + int rank = ((args->proc*args->nThreads + args->thread)*args->nGpus + i); + HIPCHECK(hipMemset(args->recvbuffs[i], 0, args->expectedBytes)); + void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i]; + if (rank == root) { + for (int j=0; jexpectedBytes*j, recvcount, type, rep, j)); + } +#ifdef DEBUG_PRINT + int *dataHost = (int *)malloc(args->sendBytes); + hipMemcpy(dataHost, data, args->sendBytes, hipMemcpyDeviceToHost); + printf("\n Rank [%d] Init: ", rank); + for (int j=0; jsendBytes/sizeof(int); j++) { + printf("%d:%d ", j, dataHost[j]); + } + printf("\n"); + free(dataHost); +#endif + } + TESTCHECK(InitData(args->expected[i], recvcount, type, rep, rank)); + HIPCHECK(hipDeviceSynchronize()); + } + return testSuccess; +} + +void ScatterGetBw(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 ScatterRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream) { + int nRanks; + NCCLCHECK(ncclCommCount(comm, &nRanks)); + size_t rankOffset = count * wordSize(type); + if (count == 0) return testSuccess; + + int rank; + NCCLCHECK(ncclCommUserRank(comm, &rank)); + NCCLCHECK(ncclGroupStart()); + if (rank == root) { + for (int r=0; rcollTest = &scatterTest; + 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 = ncclNumTypes; + run_types = test_types; + run_typenames = test_typenames; + } + + for (int i=0; i