Merge remote-tracking branch 'nccl/p2p' into p2p
This commit is contained in:
Wenkai Du
+1
-1
@@ -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
|
||||
BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall
|
||||
BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf)
|
||||
|
||||
build: ${BIN_FILES}
|
||||
|
||||
@@ -0,0 +1,116 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2020, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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",
|
||||
"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 AlltoAllGetCollByteCount(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)*nranks;
|
||||
*sendInplaceOffset = 0;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = count/nranks;
|
||||
}
|
||||
|
||||
testResult_t AlltoAllInitData(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; i<args->nGpus; i++) {
|
||||
char* str = getenv("NCCL_TESTS_DEVICE");
|
||||
int gpuid = str ? atoi(str) : 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];
|
||||
TESTCHECK(InitData(data, sendcount, type, rep, rank));
|
||||
for (int j=0; j<nranks; j++) {
|
||||
TESTCHECK(InitData(((char*)args->expected[i])+args->sendBytes/nranks*j, sendcount/nranks, type, rep+rank*sendcount/nranks, j));
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
// We don't support in-place alltoall
|
||||
args->reportErrors = in_place ? 0 : 1;
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void AlltoAllGetBw(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 AlltoAllRunColl(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;
|
||||
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
for (int r=0; r<nRanks; r++) {
|
||||
NCCLCHECK(ncclSend(((char*)sendbuff)+r*rankOffset, count, type, r, comm, stream));
|
||||
NCCLCHECK(ncclRecv(((char*)recvbuff)+r*rankOffset, count, type, r, comm, stream));
|
||||
}
|
||||
NCCLCHECK(ncclGroupEnd());
|
||||
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl alltoAllTest = {
|
||||
"AlltoAll",
|
||||
AlltoAllGetCollByteCount,
|
||||
AlltoAllInitData,
|
||||
AlltoAllGetBw,
|
||||
AlltoAllRunColl
|
||||
};
|
||||
|
||||
void AlltoAllGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
AlltoAllGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
}
|
||||
|
||||
testResult_t AlltoAllRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
args->collTest = &alltoAllTest;
|
||||
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<type_count; i++) {
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t)0, "", -1));
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = {
|
||||
AlltoAllGetBuffSize,
|
||||
AlltoAllRunTest
|
||||
};
|
||||
+4
-6
@@ -355,11 +355,7 @@ testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
|
||||
#endif
|
||||
}
|
||||
double nranks = args->nProcs*args->nThreads*args->nGpus;
|
||||
if (maxDelta > DeltaMaxValue(type)*(nranks - 1))
|
||||
{
|
||||
args->errors[0]++;
|
||||
*error = true;
|
||||
}
|
||||
if (args->reportErrors && maxDelta > DeltaMaxValue(type)*(nranks - 1)) args->errors[0]++;
|
||||
*delta = maxDelta;
|
||||
return testSuccess;
|
||||
}
|
||||
@@ -575,7 +571,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 %d cycle.\n", iter+1);
|
||||
if (iter > 0) PRINT("# Testing %ld 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);
|
||||
@@ -959,6 +955,8 @@ testResult_t run() {
|
||||
threads[t].args.bw=bw+t;
|
||||
threads[t].args.bw_count=bw_count+t;
|
||||
|
||||
threads[t].args.reportErrors = 1;
|
||||
|
||||
threads[t].func = parallel_init ? threadInit : threadRunTests;
|
||||
if (t)
|
||||
TESTCHECK(threadLaunch(threads+t));
|
||||
|
||||
+5
-3
@@ -29,13 +29,13 @@
|
||||
} while(0)
|
||||
|
||||
#define NCCLCHECK(cmd) do { \
|
||||
ncclResult_t r = cmd; \
|
||||
if (r!= ncclSuccess) { \
|
||||
ncclResult_t res = cmd; \
|
||||
if (res != ncclSuccess) { \
|
||||
char hostname[1024]; \
|
||||
getHostName(hostname, 1024); \
|
||||
printf("%s: Test NCCL failure %s:%d '%s'\n", \
|
||||
hostname, \
|
||||
__FILE__,__LINE__,ncclGetErrorString(r)); \
|
||||
__FILE__,__LINE__,ncclGetErrorString(res)); \
|
||||
return testNcclError; \
|
||||
} \
|
||||
} while(0)
|
||||
@@ -124,6 +124,8 @@ struct threadArgs {
|
||||
double* bw;
|
||||
int* bw_count;
|
||||
|
||||
int reportErrors;
|
||||
|
||||
struct testColl* collTest;
|
||||
};
|
||||
|
||||
|
||||
مرجع در شماره جدید
Block a user