/************************************************************************* * Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved. * Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved. * * See LICENSE.txt for license information ************************************************************************/ #include "cuda_runtime.h" #include "common.h" #include "rccl_compat.h" void GatherGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) { *sendcount = (count/nranks) & -(16/eltSize); *recvcount = (*sendcount)*nranks; *sendInplaceOffset = *sendcount; *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++) { CUDACHECK(cudaSetDevice(args->gpus[i])); 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, rank*sendcount, type, ncclSum, rep, 1, 0)); CUDACHECK(cudaMemcpy(args->expected[i], args->recvbuffs[i], args->expectedBytes, cudaMemcpyDefault)); if (rank == root) { TESTCHECK(InitData(args->expected[i], nranks*sendcount, 0, type, ncclSum, rep, 1, 0)); } 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, void* bias = nullptr) { 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