/************************************************************************* * Copyright (c) 2015-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 #include "common.h" void BroadcastGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) { *sendcount = count; *recvcount = count; *sendInplaceOffset = 0; *recvInplaceOffset = 0; *paramcount = *sendcount; } testResult_t BroadcastInitData(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 k=0; for (int i=0; inGpus; i++) { HIPCHECK(hipSetDevice(args->gpus[i])); for (int l=0; lnRanks; l++) { int rank = ((args->proc*args->nThreads + args->thread)*args->nGpus*args->nRanks + i*args->nRanks + l); HIPCHECK(hipMemset(args->recvbuffs[k], 0, args->expectedBytes)); void* data = in_place ? args->recvbuffs[k] : args->sendbuffs[k]; if (rank == root) TESTCHECK(InitData(data, sendcount, 0, type, ncclSum, rep, 1, 0)); TESTCHECK(InitData(args->expected[k], recvcount, 0, type, ncclSum, rep, 1, 0)); k++; } HIPCHECK(hipDeviceSynchronize()); } return testSuccess; } void BroadcastGetBw(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 BroadcastRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream) { int rank; NCCLCHECK(ncclCommUserRank(comm, &rank)); #if NCCL_MAJOR >= 2 && NCCL_MINOR >= 2 NCCLCHECK(ncclBroadcast(sendbuff, recvbuff, count, type, root, comm, stream)); #else if (rank == root) { NCCLCHECK(ncclBcast(sendbuff, count, type, root, comm, stream)); } else { NCCLCHECK(ncclBcast(recvbuff, count, type, root, comm, stream)); } #endif return testSuccess; } struct testColl broadcastTest = { "Broadcast", BroadcastGetCollByteCount, BroadcastInitData, BroadcastGetBw, BroadcastRunColl }; void BroadcastGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) { size_t paramcount, sendInplaceOffset, recvInplaceOffset; BroadcastGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks); } testResult_t BroadcastRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) { args->collTest = &broadcastTest; 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