Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl-tests commit: 9f8ddadcdf]
Dieser Commit ist enthalten in:
Wenkai Du
@@ -1,6 +1,5 @@
|
||||
#
|
||||
# Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
|
||||
# Modifications are Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# See LICENCE.txt for license information
|
||||
#
|
||||
|
||||
@@ -52,17 +52,22 @@ All tests support the same set of arguments :
|
||||
* `-i,--stepbytes <increment size>` fixed increment between sizes. Default : (max-min)/10.
|
||||
* `-f,--stepfactor <increment factor>` multiplication factor between sizes. Default : disabled.
|
||||
* RCCL operations arguments
|
||||
* `-o,--op <sum/prod/min/max/all>` Specify which reduction operation to perform. Only relevant for reduction operations like Allreduce, Reduce or ReduceScatter. Default : Sum.
|
||||
* `-o,--op <sum/prod/min/max/avg/all>` Specify which reduction operation to perform. Only relevant for reduction operations like Allreduce, Reduce or ReduceScatter. Default : Sum.
|
||||
* `-d,--datatype <nccltype/all>` Specify which datatype to use. Default : Float.
|
||||
* `-r,--root <root/all>` Specify which root to use. Only for operations with a root like broadcast or reduce. Default : 0.
|
||||
* `-y,--memory_type <coarse/fine/host/managed>` Default: Coarse
|
||||
* `-s,--stress_cycles <number of cycles>` Default: 1
|
||||
* `-u,--cumask <d0,d1,d2,d3>` Default: None
|
||||
* Performance
|
||||
* `-n,--iters <iteration count>` number of iterations. Default : 20.
|
||||
* `-w,--warmup_iters <warmup iteration count>` number of warmup iterations (not timed). Default : 5.
|
||||
* `-m,--agg_iters <aggregation count>` 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 RCCL 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 RCCL collective blocking, i.e. have CPUs wait and sync after each collective. Default : 0.
|
||||
* `-G,--cudagraph <num graph launches>` Capture iterations as a CUDA graph and then replay specified number of times. Default : 0.
|
||||
|
||||
## Unit tests
|
||||
|
||||
@@ -80,7 +85,7 @@ $ LD_LIBRARY_PATH=/path/to/rccl-install/lib/ HSA_FORCE_FINE_GRAIN_PCIE=1 python3
|
||||
|
||||
RCCL tests are provided under the BSD license.
|
||||
|
||||
All source code and accompanying documentation is copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved.
|
||||
All source code and accompanying documentation is copyright (c) 2016-2021, NVIDIA CORPORATION. All rights reserved.
|
||||
|
||||
All modifications are copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
|
||||
@@ -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 RCCL 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`
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
#
|
||||
# Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
|
||||
# Copyright (c) 2015-2021, NVIDIA CORPORATION. All rights reserved.
|
||||
# Modifications are Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# See LICENSE.txt for license information
|
||||
@@ -57,7 +57,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 gather scatter sendrecv alltoallv
|
||||
BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall gather scatter sendrecv alltoallv hypercube
|
||||
BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf)
|
||||
|
||||
build: ${BIN_FILES}
|
||||
|
||||
@@ -9,15 +9,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) {
|
||||
@@ -85,7 +85,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;
|
||||
}
|
||||
|
||||
@@ -9,15 +9,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) {
|
||||
@@ -84,7 +84,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;
|
||||
}
|
||||
@@ -94,7 +94,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;
|
||||
}
|
||||
|
||||
@@ -8,18 +8,16 @@
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "common.h"
|
||||
|
||||
#define USE_RCCL_GATHER_SCATTER
|
||||
|
||||
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) {
|
||||
@@ -70,9 +68,6 @@ testResult_t AlltoAllRunColl(void* sendbuff, void* recvbuff, size_t count, ncclD
|
||||
#if NCCL_MAJOR < 2 || NCCL_MINOR < 7
|
||||
printf("NCCL 2.7 or later is needed for alltoall. This test was compiled with %d.%d.\n", NCCL_MAJOR, NCCL_MINOR);
|
||||
return testNcclError;
|
||||
#else
|
||||
#if defined(RCCL_GATHER_SCATTER) && defined(USE_RCCL_GATHER_SCATTER)
|
||||
NCCLCHECK(ncclAllToAll(sendbuff, recvbuff, count, type, comm, stream));
|
||||
#else
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
for (int r=0; r<nRanks; r++) {
|
||||
@@ -80,7 +75,6 @@ testResult_t AlltoAllRunColl(void* sendbuff, void* recvbuff, size_t count, ncclD
|
||||
NCCLCHECK(ncclRecv(((char*)recvbuff)+r*rankOffset, count, type, r, comm, stream));
|
||||
}
|
||||
NCCLCHECK(ncclGroupEnd());
|
||||
#endif
|
||||
return testSuccess;
|
||||
#endif
|
||||
}
|
||||
@@ -109,7 +103,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;
|
||||
}
|
||||
|
||||
@@ -9,15 +9,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) {
|
||||
@@ -93,7 +93,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;
|
||||
}
|
||||
|
||||
@@ -1,7 +1,7 @@
|
||||
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
@@ -17,20 +17,35 @@
|
||||
|
||||
//#define DEBUG_PRINT
|
||||
|
||||
int test_ncclVersion = 0; // init'd with ncclGetVersion()
|
||||
|
||||
#if NCCL_MAJOR >= 2
|
||||
#if RCCL_BFLOAT16 == 1
|
||||
ncclDataType_t test_types[ncclNumTypes] = {ncclInt8, ncclUint8, ncclInt32, ncclUint32, ncclInt64, ncclUint64, ncclHalf, ncclFloat, ncclDouble, ncclBfloat16};
|
||||
const char *test_typenames[ncclNumTypes] = {"int8", "uint8", "int32", "uint32", "int64", "uint64", "half", "float", "double", "bf16"};
|
||||
int test_typenum = 10;
|
||||
#else
|
||||
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"};
|
||||
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
|
||||
|
||||
const char *test_memorytypes[nccl_NUM_MTYPES] = {"coarse", "fine", "host", "managed"};
|
||||
|
||||
thread_local int is_main_thread = 0;
|
||||
@@ -54,28 +69,54 @@ static int blocking_coll = 0;
|
||||
static int memorytype = 0;
|
||||
static int stress_cycles = 1;
|
||||
static uint32_t cumask[4];
|
||||
static int cudaGraphLaunches = 0;
|
||||
// Report average iteration time: (0=RANK0,1=AVG,2=MIN,3=MAX)
|
||||
static int average = 1;
|
||||
|
||||
double parsesize(char *value) {
|
||||
#define NUM_BLOCKS 32
|
||||
|
||||
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) {
|
||||
switch(type) {
|
||||
case ncclHalf: return 1e-2;
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
case ncclBfloat16: return 1e-2;
|
||||
#endif
|
||||
case ncclFloat: return 1e-5;
|
||||
case ncclDouble: return 1e-12;
|
||||
case ncclInt:
|
||||
@@ -86,9 +127,6 @@ double DeltaMaxValue(ncclDataType_t type) {
|
||||
#endif
|
||||
case ncclInt64:
|
||||
case ncclUint64: return 1e-200;
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
case ncclBfloat16: return 1e-2;
|
||||
#endif
|
||||
}
|
||||
return 1e-200;
|
||||
}
|
||||
@@ -113,25 +151,32 @@ 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<typename T, int BSIZE> __global__
|
||||
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<count; i+=blockDim.x) {
|
||||
for(size_t i=tid; i<count; i+=blockDim.x*gridDim.x) {
|
||||
|
||||
double delta = absDiff(A[i], B[i]);
|
||||
if( delta > locmax ) {
|
||||
locmax = delta;
|
||||
#if 0
|
||||
if (delta > .1) printf("Error at %d/%ld : %f != %f\n", i, count, toFloat(A[i]), toFloat(B[i]));
|
||||
#ifdef DEBUG_PRINT
|
||||
if (delta > .1) printf("Error at %ld/%ld(%p) : %f != %f\n", i, count, B+i, toFloat(A[i]), toFloat(B[i]));
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
tid = threadIdx.x;
|
||||
temp[tid] = locmax;
|
||||
for(int stride = BSIZE/2; stride > 1; stride>>=1) {
|
||||
__syncthreads();
|
||||
@@ -140,19 +185,21 @@ 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 NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
case ncclBfloat16:
|
||||
hipLaunchKernelGGL((deltaKern<rccl_bfloat16, 512>), dim3(1), dim3(512), 0, 0, results, expected, count, devmax); break;
|
||||
#endif
|
||||
case ncclHalf:
|
||||
hipLaunchKernelGGL((deltaKern<half, 512>), dim3(1), dim3(512), 0, 0, results, expected, count, devmax); break;
|
||||
case ncclFloat:
|
||||
hipLaunchKernelGGL((deltaKern<float, 512>), dim3(1), dim3(512), 0, 0, results, expected, count, devmax); break;
|
||||
case ncclDouble:
|
||||
hipLaunchKernelGGL((deltaKern<double, 512>), dim3(1), dim3(512), 0, 0, results, expected, count, devmax); break;
|
||||
|
||||
case ncclChar:
|
||||
#if NCCL_MAJOR >= 2
|
||||
case ncclUint8:
|
||||
@@ -166,12 +213,9 @@ testResult_t CheckDelta(void* expected, void* results, size_t count, ncclDataTyp
|
||||
case ncclInt64:
|
||||
case ncclUint64:
|
||||
hipLaunchKernelGGL((deltaKern<uint64_t, 512>), dim3(1), dim3(512), 0, 0, results, expected, count, devmax); break;
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
case ncclBfloat16:
|
||||
hipLaunchKernelGGL((deltaKern<rccl_bfloat16, 512>), dim3(1), dim3(512), 0, 0, results, expected, count, devmax); break;
|
||||
#endif
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
for (int i=1; i<NUM_BLOCKS; i++) devmax[0] = std::max(devmax[0], devmax[i]);
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
@@ -196,10 +240,12 @@ template<>
|
||||
__device__ half testValue<half>(const size_t offset, const int rep, const int rank) {
|
||||
return __float2half(testValue<float>(offset, rep, rank));
|
||||
}
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
template<>
|
||||
__device__ rccl_bfloat16 testValue<rccl_bfloat16>(const size_t offset, const int rep, const int rank) {
|
||||
return rccl_bfloat16(testValue<float>(offset, rep, rank));
|
||||
}
|
||||
#endif
|
||||
|
||||
// Operations
|
||||
template<typename T>
|
||||
@@ -221,71 +267,78 @@ __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<typename T, T (*Op)(T, T)>
|
||||
__global__ void InitDataReduceKernel(void* data, const size_t N, const size_t offset, const int rep, const int nranks) {
|
||||
template<typename T>
|
||||
__device__ T ncclPostOpIdent(T x, int n) { return x; }
|
||||
|
||||
template<typename T>
|
||||
__device__ T ncclPostOpDiv(T x, int n) { return x/n; }
|
||||
template<>
|
||||
__device__ half ncclPostOpDiv<half>(half x, int n) { return __float2half(__half2float(x)/n); }
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
template<>
|
||||
__device__ rccl_bfloat16 ncclPostOpDiv<rccl_bfloat16>(rccl_bfloat16 x, int n) { return (rccl_bfloat16)((float)x/n); }
|
||||
#endif
|
||||
|
||||
template<typename T, T (*Op)(T, T), T(*PostOp)(T,int)>
|
||||
__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<N; o+=gridDim.x*blockDim.x) {
|
||||
T val = testValue<T>(o+offset, rep, 0);
|
||||
for (int i=1; i<nranks; i++) {
|
||||
val = Op(val, testValue<T>(o+offset, rep, i));
|
||||
}
|
||||
((T*)data)[o] = val;
|
||||
data[o] = PostOp(val, nranks);
|
||||
}
|
||||
}
|
||||
|
||||
typedef void(*redInitKern_t)(void* data, const size_t N, const size_t offset, const int rep, const int nranks);
|
||||
|
||||
#define KERN(type, op) InitDataReduceKernel<type, op<type>>
|
||||
#define OPS(type) KERN(type, ncclOpSum), KERN(type, ncclOpProd), KERN(type, ncclOpMax), KERN(type, ncclOpMin)
|
||||
|
||||
static redInitKern_t const redInitDataKerns[ncclNumOps*ncclNumTypes] = {
|
||||
#if NCCL_MAJOR >= 2
|
||||
#if RCCL_BFLOAT16 == 1
|
||||
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(rccl_bfloat16)
|
||||
#define KERN(type, op, postop) (void*)InitDataReduceKernel<type, op<type>, postop<type> >
|
||||
#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
|
||||
OPS(int8_t), OPS(uint8_t), OPS(int32_t), OPS(uint32_t), OPS(int64_t), OPS(uint64_t), OPS(half), OPS(float), OPS(double)
|
||||
#define OPS(type) \
|
||||
KERN(type, ncclOpSum, ncclPostOpIdent), \
|
||||
KERN(type, ncclOpProd, ncclPostOpIdent), \
|
||||
KERN(type, ncclOpMax, ncclPostOpIdent), \
|
||||
KERN(type, ncclOpMin, ncclPostOpIdent)
|
||||
#endif
|
||||
#else
|
||||
OPS(char), OPS(int32_t), OPS(half), OPS(float), OPS(double), OPS(int64_t), OPS(uint64_t)
|
||||
|
||||
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 NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
OPS(rccl_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) {
|
||||
dim3 grid = { 32, 1, 1 };
|
||||
dim3 block = { 256, 1, 1 };
|
||||
hipLaunchKernelGGL((redInitDataKerns[type*ncclNumOps+op]), grid, block, 0, 0, data, count, offset, rep, nranks);
|
||||
void* args[5] = { (void*)&data, (void*)&count, (void*)&offset, (void*)&rep, (void*)&nranks };
|
||||
HIPCHECK(hipLaunchKernel(redInitDataKerns[type*ncclNumOps+op], grid, block, args, 0, hipStreamDefault));
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
__global__ void InitDataKernel(void* data, const size_t N, const int rep, const int rank) {
|
||||
__global__ void InitDataKernel(T* data, const size_t N, const int rep, const int rank) {
|
||||
for (size_t o=blockIdx.x*blockDim.x+threadIdx.x; o<N; o+=gridDim.x*blockDim.x)
|
||||
((T*)data)[o] = testValue<T>(o, rep, rank);
|
||||
data[o] = testValue<T>(o, rep, rank);
|
||||
}
|
||||
|
||||
typedef void(*initDataKern_t)(void* data, const size_t N, const int rep, const int rank);
|
||||
|
||||
static initDataKern_t const initDataKerns[ncclNumTypes] = {
|
||||
#if NCCL_MAJOR >= 2
|
||||
InitDataKernel< int8_t>,
|
||||
InitDataKernel< uint8_t>,
|
||||
InitDataKernel< int32_t>,
|
||||
InitDataKernel<uint32_t>,
|
||||
InitDataKernel< int64_t>,
|
||||
InitDataKernel<uint64_t>,
|
||||
InitDataKernel< half>,
|
||||
InitDataKernel< float>,
|
||||
InitDataKernel< double>,
|
||||
#if RCCL_BFLOAT16 == 1
|
||||
InitDataKernel<rccl_bfloat16>
|
||||
#endif
|
||||
#else
|
||||
InitDataKernel< char>,
|
||||
InitDataKernel< int32_t>,
|
||||
InitDataKernel< half>,
|
||||
InitDataKernel< float>,
|
||||
InitDataKernel< double>,
|
||||
InitDataKernel< int64_t>,
|
||||
InitDataKernel<uint64_t>,
|
||||
static void* const initDataKerns[ncclNumTypes] = {
|
||||
(void*)InitDataKernel< int8_t>,
|
||||
(void*)InitDataKernel< uint8_t>,
|
||||
(void*)InitDataKernel< int32_t>,
|
||||
(void*)InitDataKernel<uint32_t>,
|
||||
(void*)InitDataKernel< int64_t>,
|
||||
(void*)InitDataKernel<uint64_t>,
|
||||
(void*)InitDataKernel< half>,
|
||||
(void*)InitDataKernel< float>,
|
||||
(void*)InitDataKernel< double>,
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
(void*)InitDataKernel<rccl_bfloat16>,
|
||||
#endif
|
||||
};
|
||||
|
||||
@@ -299,16 +352,14 @@ testResult_t InitDataType(void* dest, const size_t N, const int rep, const int r
|
||||
testResult_t InitData(void* data, const size_t count, ncclDataType_t type, const int rep, const int rank) {
|
||||
dim3 grid = { 32, 1, 1 };
|
||||
dim3 block = { 256, 1, 1 };
|
||||
hipLaunchKernelGGL((initDataKerns[type]), grid, block, 0, 0, data, count, rep, rank);
|
||||
void* args[4] = { (void*)&data, (void*)&count, (void*)&rep, (void*)&rank };
|
||||
HIPCHECK(hipLaunchKernel(initDataKerns[type], grid, block, args, 0, hipStreamDefault));
|
||||
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);
|
||||
@@ -317,11 +368,39 @@ void Barrier(struct threadArgs* args)
|
||||
} else {
|
||||
while (args->barrier[args->barrier_idx]) pthread_yield();
|
||||
}
|
||||
|
||||
args->barrier_idx=!args->barrier_idx;
|
||||
}
|
||||
|
||||
testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double *delta, bool *error) {
|
||||
// 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;
|
||||
}
|
||||
|
||||
testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double *delta) {
|
||||
size_t count = args->expectedBytes/wordSize(type);
|
||||
double maxDelta = 0.0;
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
@@ -332,7 +411,7 @@ testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
|
||||
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];
|
||||
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
|
||||
@@ -389,9 +468,8 @@ testResult_t testStreamSynchronize(int ngpus, hipStream_t* streams, ncclComm_t*
|
||||
|
||||
if (hipErr != hipErrorNotReady) HIPCHECK(hipErr);
|
||||
|
||||
#if NCCL_MAJOR >= 2
|
||||
#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) {
|
||||
@@ -403,7 +481,6 @@ testResult_t testStreamSynchronize(int ngpus, hipStream_t* streams, ncclComm_t*
|
||||
NCCLCHECK(ncclAsyncErr);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
}
|
||||
|
||||
@@ -419,8 +496,8 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
|
||||
|
||||
// Try to change offset for each iteration so that we avoid cache effects and catch race conditions in ptrExchange
|
||||
size_t totalnbytes = max(args->sendBytes, 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++) {
|
||||
@@ -456,6 +533,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));
|
||||
@@ -463,6 +544,15 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
|
||||
|
||||
Barrier(args);
|
||||
|
||||
hipGraph_t graphs[args->nGpus];
|
||||
hipGraphExec_t graphExec[args->nGpus];
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
// Begin cuda graph capture
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipStreamBeginCapture(args->streams[i], args->nThreads > 1 ? hipStreamCaptureModeThreadLocal : hipStreamCaptureModeGlobal));
|
||||
}
|
||||
}
|
||||
|
||||
// Performance Benchmark
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
for (int iter = 0; iter < iters; iter++) {
|
||||
@@ -472,11 +562,41 @@ 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; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipStreamEndCapture(args->streams[i], graphs+i));
|
||||
}
|
||||
// Instantiate cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipGraphInstantiate(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; l<cudaGraphLaunches; l++) {
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipGraphLaunch(graphExec[i], args->streams[i]));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
TESTCHECK(completeColl(args));
|
||||
|
||||
auto delta = std::chrono::high_resolution_clock::now() - start;
|
||||
double deltaSec = std::chrono::duration_cast<std::chrono::duration<double>>(delta).count();
|
||||
deltaSec = deltaSec/(iters*agg_iters);
|
||||
if (cudaGraphLaunches >= 1) deltaSec = deltaSec/cudaGraphLaunches;
|
||||
Allreduce(args, &deltaSec, average);
|
||||
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
//destroy cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipGraphExecDestroy(graphExec[i]));
|
||||
HIPCHECK(hipGraphDestroy(graphs[i]));
|
||||
}
|
||||
}
|
||||
|
||||
double algBw, busBw;
|
||||
args->collTest->getBw(count, wordSize(type), deltaSec, &algBw, &busBw, args->nProcs*args->nThreads*args->nGpus);
|
||||
@@ -491,31 +611,52 @@ 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; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipStreamBeginCapture(args->streams[i], hipStreamCaptureModeThreadLocal));
|
||||
}
|
||||
}
|
||||
|
||||
//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; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipStreamEndCapture(args->streams[i], graphs+i));
|
||||
}
|
||||
// Instantiate cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0));
|
||||
}
|
||||
// Launch cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipGraphLaunch(graphExec[i], args->streams[i]));
|
||||
}
|
||||
}
|
||||
|
||||
TESTCHECK(completeColl(args));
|
||||
|
||||
TESTCHECK(CheckData(args, type, op, root, in_place, &maxDelta, &error));
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
//destroy cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
HIPCHECK(hipGraphExecDestroy(graphExec[i]));
|
||||
HIPCHECK(hipGraphDestroy(graphs[i]));
|
||||
}
|
||||
}
|
||||
|
||||
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; i<args->nThreads; i++) {
|
||||
maxDelta += args->deltaThreads[i];
|
||||
}
|
||||
#ifdef MPI_SUPPORT
|
||||
MPI_Allreduce(MPI_IN_PLACE, &maxDelta, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD);
|
||||
MPI_Allreduce(MPI_IN_PLACE, &error, 1, MPI_C_BOOL, MPI_LOR, MPI_COMM_WORLD);
|
||||
#endif
|
||||
}
|
||||
Barrier(args);
|
||||
Allreduce(args, &maxDelta, 3);
|
||||
}
|
||||
|
||||
double timeUsec = deltaSec*1.0E6;
|
||||
char timeStr[10];
|
||||
if (timeUsec > 10000.0) {
|
||||
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);
|
||||
@@ -604,11 +745,7 @@ testResult_t threadInit(struct threadArgs* args) {
|
||||
TESTCHECK(threadRunTests(args));
|
||||
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
#if NCCL_MAJOR >= 2
|
||||
NCCLCHECK(ncclCommDestroy(args->comms[i]));
|
||||
#else
|
||||
ncclCommDestroy(args->comms[i]);
|
||||
#endif
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
@@ -627,27 +764,27 @@ testResult_t AllocateBuffs(void **sendbuff, size_t sendBytes, void **recvbuff, s
|
||||
if (memorytype == ncclFine) {
|
||||
HIPCHECK(hipExtMallocWithFlags(sendbuff, nbytes, hipDeviceMallocFinegrained));
|
||||
HIPCHECK(hipExtMallocWithFlags(recvbuff, nbytes, hipDeviceMallocFinegrained));
|
||||
HIPCHECK(hipExtMallocWithFlags(expected, recvBytes, hipDeviceMallocFinegrained));
|
||||
if (datacheck) HIPCHECK(hipExtMallocWithFlags(expected, recvBytes, hipDeviceMallocFinegrained));
|
||||
}
|
||||
else if (memorytype == ncclHost) {
|
||||
HIPCHECK(hipHostMalloc(sendbuff, nbytes));
|
||||
HIPCHECK(hipHostMalloc(recvbuff, nbytes));
|
||||
HIPCHECK(hipHostMalloc(expected, recvBytes));
|
||||
if (datacheck) HIPCHECK(hipHostMalloc(expected, recvBytes));
|
||||
}
|
||||
else if (memorytype == ncclManaged) {
|
||||
HIPCHECK(hipMallocManaged(sendbuff, nbytes));
|
||||
HIPCHECK(hipMallocManaged(recvbuff, nbytes));
|
||||
HIPCHECK(hipMallocManaged(expected, recvBytes));
|
||||
if (datacheck) HIPCHECK(hipMallocManaged(expected, recvBytes));
|
||||
#if 0
|
||||
HIPCHECK(hipMemset(*sendbuff, 0, nbytes));
|
||||
HIPCHECK(hipMemset(*recvbuff, 0, nbytes));
|
||||
HIPCHECK(hipMemset(*expected, 0, recvBytes));
|
||||
if (datacheck) HIPCHECK(hipMemset(*expected, 0, recvBytes));
|
||||
#endif
|
||||
}
|
||||
else {
|
||||
HIPCHECK(hipMalloc(sendbuff, nbytes));
|
||||
HIPCHECK(hipMalloc(recvbuff, nbytes));
|
||||
HIPCHECK(hipMalloc(expected, recvBytes));
|
||||
if (datacheck) HIPCHECK(hipMalloc(expected, recvBytes));
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
@@ -658,7 +795,19 @@ 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
|
||||
double parsed;
|
||||
int longindex;
|
||||
static struct option longopts[] = {
|
||||
{"nthreads", required_argument, 0, 't'},
|
||||
@@ -679,12 +828,14 @@ int main(int argc, char* argv[]) {
|
||||
{"memory_type", required_argument, 0, 'y'},
|
||||
{"stress_cycles", required_argument, 0, 's'},
|
||||
{"cumask", required_argument, 0, 'u'},
|
||||
{"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:y:s:u:h", longopts, &longindex);
|
||||
c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:p:c:o:d:r:z:G:a:y:s:u:h:", longopts, &longindex);
|
||||
|
||||
if (c == -1)
|
||||
break;
|
||||
@@ -697,10 +848,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);
|
||||
@@ -712,10 +873,10 @@ 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");
|
||||
fprintf(stderr, "Option -m not supported before NCCL 2.2. Ignoring\n");
|
||||
#endif
|
||||
break;
|
||||
case 'w':
|
||||
@@ -755,32 +916,20 @@ int main(int argc, char* argv[]) {
|
||||
};
|
||||
}
|
||||
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 'a':
|
||||
average = (int)strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'h':
|
||||
printf("USAGE: %s \n\t"
|
||||
"[-t,--nthreads <num threads>] \n\t"
|
||||
"[-g,--ngpus <gpus per thread>] \n\t"
|
||||
"[-b,--minbytes <min size in bytes>] \n\t"
|
||||
"[-e,--maxbytes <max size in bytes>] \n\t"
|
||||
"[-i,--stepbytes <increment size>] \n\t"
|
||||
"[-f,--stepfactor <increment factor>] \n\t"
|
||||
"[-n,--iters <iteration count>] \n\t"
|
||||
"[-m,--agg_iters <aggregated iteration count>] \n\t"
|
||||
"[-w,--warmup_iters <warmup iteration count>] \n\t"
|
||||
"[-p,--parallel_init <0/1>] \n\t"
|
||||
"[-c,--check <0/1>] \n\t"
|
||||
"[-o,--op <sum/prod/min/max/all>] \n\t"
|
||||
"[-d,--datatype <nccltype/all>] \n\t"
|
||||
"[-r,--root <root>] \n\t"
|
||||
"[-z,--blocking <0/1>] \n\t"
|
||||
"[-y,--memory_type <coarse/fine/host>] \n\t"
|
||||
"[-s,--stress_cycles <number of cycles>] \n\t"
|
||||
"[-u,--cumask <d0,d1,d2,d3>] \n\t"
|
||||
"[-h,--help]\n",
|
||||
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 <num threads>] \n\t"
|
||||
"[-g,--ngpus <gpus per thread>] \n\t"
|
||||
"[-b,--minbytes <min size in bytes>] \n\t"
|
||||
@@ -792,13 +941,19 @@ int main(int argc, char* argv[]) {
|
||||
"[-w,--warmup_iters <warmup iteration count>] \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 <sum/prod/min/max/avg/all>] \n\t"
|
||||
#else
|
||||
"[-o,--op <sum/prod/min/max/all>] \n\t"
|
||||
#endif
|
||||
"[-d,--datatype <nccltype/all>] \n\t"
|
||||
"[-r,--root <root>] \n\t"
|
||||
"[-z,--blocking <0/1>] \n\t"
|
||||
"[-y,--memory_type <coarse/fine/host/managed>] \n\t"
|
||||
"[-s,--stress_cycles <number of cycles>] \n\t"
|
||||
"[-u,--cumask <d0,d1,d2,d3>] \n\t"
|
||||
"[-G,--cudagraph <num graph launches>] \n\t"
|
||||
"[-a,--average <0/1/2/3> report average iteration time <0=RANK0/1=AVG/2=MIN/3=MAX>] \n\t"
|
||||
"[-h,--help]\n",
|
||||
basename(argv[0]));
|
||||
return 0;
|
||||
@@ -813,10 +968,17 @@ int main(int argc, char* argv[]) {
|
||||
return testNcclError;
|
||||
}
|
||||
|
||||
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
|
||||
return run();
|
||||
TESTCHECK(run());
|
||||
return 0;
|
||||
}
|
||||
|
||||
testResult_t run() {
|
||||
@@ -848,6 +1010,7 @@ testResult_t run() {
|
||||
#define MAX_LINE 2048
|
||||
char line[MAX_LINE];
|
||||
int len = 0;
|
||||
size_t maxMem = ~0;
|
||||
for (int i=0; i<nThreads*nGpus; i++) {
|
||||
int hipDev = localRank*nThreads*nGpus+i;
|
||||
int rank = proc*nThreads*nGpus+i;
|
||||
@@ -855,6 +1018,7 @@ testResult_t run() {
|
||||
HIPCHECK(hipGetDeviceProperties(&prop, hipDev));
|
||||
len += snprintf(line+len, MAX_LINE>len ? MAX_LINE-len : 0, "# Rank %2d Pid %6d on %10s device %2d [0x%02x] %s\n",
|
||||
rank, getpid(), hostname, hipDev, prop.pciBusID, prop.name);
|
||||
maxMem = std::min(maxMem, prop.totalGlobalMem);
|
||||
}
|
||||
|
||||
#if MPI_SUPPORT
|
||||
@@ -866,10 +1030,18 @@ testResult_t run() {
|
||||
PRINT("%s", lines+MAX_LINE*p);
|
||||
free(lines);
|
||||
}
|
||||
MPI_Allreduce(MPI_IN_PLACE, &maxMem, 1, MPI_LONG, MPI_MIN, MPI_COMM_WORLD);
|
||||
#else
|
||||
PRINT("%s", line);
|
||||
#endif
|
||||
|
||||
// We need sendbuff, recvbuff, expected (when datacheck enabled), plus 1G for the rest.
|
||||
size_t memMaxBytes = (maxMem - (1<<30)) / (datacheck ? 3 : 2);
|
||||
if (maxBytes > 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));
|
||||
@@ -897,11 +1069,7 @@ testResult_t run() {
|
||||
} else
|
||||
HIPCHECK(hipStreamCreateWithFlags(streams+i, hipStreamNonBlocking));
|
||||
// initialize data buffer to avoid all zero data
|
||||
#if NCCL_MAJOR >= 2
|
||||
TESTCHECK(InitData(sendbuffs[i], sendBytes, ncclUint8, 0, i));
|
||||
#else
|
||||
TESTCHECK(InitData(sendbuffs[i], sendBytes, ncclChar, 0, i));
|
||||
#endif
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
@@ -925,7 +1093,7 @@ testResult_t run() {
|
||||
int errors[nThreads];
|
||||
double bw[nThreads];
|
||||
double* delta;
|
||||
HIPCHECK(hipHostMalloc(&delta, sizeof(double)*nThreads, hipHostMallocPortable | hipHostMallocMapped));
|
||||
HIPCHECK(hipHostMalloc(&delta, sizeof(double)*nThreads*NUM_BLOCKS, hipHostMallocPortable | hipHostMallocMapped));
|
||||
int bw_count[nThreads];
|
||||
for (int t=0; t<nThreads; t++) {
|
||||
bw[t] = 0.0;
|
||||
@@ -937,6 +1105,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);
|
||||
@@ -962,11 +1131,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);
|
||||
threads[t].args.delta = delta;
|
||||
threads[t].args.deltaHost = (delta + t*NUM_BLOCKS);
|
||||
threads[t].args.errors=errors+t;
|
||||
threads[t].args.bw=bw+t;
|
||||
threads[t].args.bw_count=bw_count+t;
|
||||
@@ -984,7 +1152,6 @@ testResult_t run() {
|
||||
for (int t=nThreads-1; t>=0; t--) {
|
||||
if (t) pthread_join(threads[t].thread, NULL);
|
||||
TESTCHECK(threads[t].ret);
|
||||
|
||||
if (t) {
|
||||
errors[0] += errors[t];
|
||||
bw[0] += bw[t];
|
||||
@@ -998,11 +1165,7 @@ testResult_t run() {
|
||||
|
||||
if (!parallel_init) {
|
||||
for(int i=0; i<nGpus*nThreads; ++i)
|
||||
#if NCCL_MAJOR >= 2
|
||||
NCCLCHECK(ncclCommDestroy(comms[i]));
|
||||
#else
|
||||
ncclCommDestroy(comms[i]);
|
||||
#endif
|
||||
free(comms);
|
||||
}
|
||||
|
||||
@@ -1011,12 +1174,12 @@ testResult_t run() {
|
||||
if (memorytype == ncclHost) {
|
||||
HIPCHECK(hipHostFree(sendbuffs[i]));
|
||||
HIPCHECK(hipHostFree(recvbuffs[i]));
|
||||
HIPCHECK(hipHostFree(expected[i]));
|
||||
if (datacheck) HIPCHECK(hipHostFree(expected[i]));
|
||||
}
|
||||
else {
|
||||
HIPCHECK(hipFree(sendbuffs[i]));
|
||||
HIPCHECK(hipFree(recvbuffs[i]));
|
||||
HIPCHECK(hipFree(expected[i]));
|
||||
if (datacheck) HIPCHECK(hipFree(expected[i]));
|
||||
}
|
||||
}
|
||||
HIPCHECK(hipHostFree(delta));
|
||||
|
||||
@@ -9,6 +9,7 @@
|
||||
|
||||
#include "rccl.h"
|
||||
#include <stdio.h>
|
||||
#include <cstdint>
|
||||
#include <algorithm>
|
||||
#ifdef MPI_SUPPORT
|
||||
#include "mpi.h"
|
||||
@@ -54,8 +55,8 @@ typedef enum {
|
||||
if (r!= testSuccess) { \
|
||||
char hostname[1024]; \
|
||||
getHostName(hostname, 1024); \
|
||||
printf(" .. %s: Test failure %s:%d\n", \
|
||||
hostname, \
|
||||
printf(" .. %s pid %d: Test failure %s:%d\n", \
|
||||
hostname, getpid(), \
|
||||
__FILE__,__LINE__); \
|
||||
return r; \
|
||||
} \
|
||||
@@ -78,6 +79,7 @@ extern struct testColl allGatherTest;
|
||||
extern struct testColl reduceScatterTest;
|
||||
extern struct testColl broadcastTest;
|
||||
extern struct testColl reduceTest;
|
||||
extern struct testColl alltoAllTest;
|
||||
|
||||
struct testEngine {
|
||||
void (*getBuffSize)(size_t *sendcount, size_t *recvcount, size_t count, int nranks);
|
||||
@@ -115,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;
|
||||
@@ -213,6 +214,9 @@ static size_t wordSize(ncclDataType_t type) {
|
||||
#endif
|
||||
return 1;
|
||||
case ncclHalf:
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
case ncclBfloat16:
|
||||
#endif
|
||||
//case ncclFloat16:
|
||||
return 2;
|
||||
case ncclInt:
|
||||
@@ -226,16 +230,13 @@ static size_t wordSize(ncclDataType_t type) {
|
||||
case ncclInt64:
|
||||
case ncclUint64:
|
||||
case ncclDouble:
|
||||
//case ncclFloat64:
|
||||
//case ncclFloat64:
|
||||
return 8;
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
case ncclBfloat16:
|
||||
return 2;
|
||||
#endif
|
||||
default: return 0;
|
||||
}
|
||||
}
|
||||
|
||||
extern int test_ncclVersion; // init'd with ncclGetVersion()
|
||||
extern ncclDataType_t test_types[ncclNumTypes];
|
||||
extern const char *test_typenames[ncclNumTypes];
|
||||
extern ncclRedOp_t test_ops[ncclNumOps];
|
||||
@@ -246,6 +247,8 @@ typedef enum { ncclCoarse = 0,
|
||||
ncclManaged = 3,
|
||||
nccl_NUM_MTYPES = 4 } ncclMemoryType_t;
|
||||
extern const char *test_memorytypes[nccl_NUM_MTYPES];
|
||||
extern int test_opnum;
|
||||
extern int test_typenum;
|
||||
|
||||
static int ncclstringtotype(char *str) {
|
||||
for (int t=0; t<ncclNumTypes; t++) {
|
||||
@@ -261,7 +264,7 @@ static int ncclstringtotype(char *str) {
|
||||
}
|
||||
|
||||
static int ncclstringtoop (char *str) {
|
||||
for (int o=0; o<ncclNumOps; o++) {
|
||||
for (int o=0; o<test_opnum; o++) {
|
||||
if (strcmp(str, test_opnames[o]) == 0) {
|
||||
return o;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,122 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2021, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "hip/hip_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; i<args->nGpus; 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));
|
||||
for (int j=0; j<nranks; j++) {
|
||||
TESTCHECK(InitData(((char*)args->expected[i])+args->sendBytes*j, sendcount, type, rep, j));
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
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, hipStream_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) HIPCHECK(hipMemcpyAsync(rbuff+rank*rankSize, sbuff, rankSize, hipMemcpyDeviceToDevice, stream));
|
||||
|
||||
// Hypercube AllGather
|
||||
for (int mask=1; mask<nRanks; mask<<=1) {
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
int s = rank & ~(mask-1);
|
||||
int r = s ^ mask;
|
||||
NCCLCHECK(ncclSend(rbuff+s*rankSize, count*mask, type, rank^mask, comm, stream));
|
||||
NCCLCHECK(ncclRecv(rbuff+r*rankSize, count*mask, type, rank^mask, comm, stream));
|
||||
NCCLCHECK(ncclGroupEnd());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl hyperCubeTest = {
|
||||
"HyperCube",
|
||||
HyperCubeGetCollByteCount,
|
||||
HyperCubeInitData,
|
||||
HyperCubeGetBw,
|
||||
HyperCubeRunColl
|
||||
};
|
||||
|
||||
void HyperCubeGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
HyperCubeGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
}
|
||||
|
||||
testResult_t HyperCubeRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
args->collTest = &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; i<type_count; i++) {
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t)0, "", -1));
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = {
|
||||
HyperCubeGetBuffSize,
|
||||
HyperCubeRunTest
|
||||
};
|
||||
@@ -9,15 +9,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) {
|
||||
@@ -84,7 +84,7 @@ testResult_t ReduceRunTest(struct threadArgs* args, int root, ncclDataType_t typ
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
} else {
|
||||
type_count = ncclNumTypes;
|
||||
type_count = test_typenum;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
@@ -94,7 +94,7 @@ testResult_t ReduceRunTest(struct threadArgs* args, int root, ncclDataType_t typ
|
||||
run_ops = &op;
|
||||
run_opnames = &opName;
|
||||
} else {
|
||||
op_count = ncclNumOps;
|
||||
op_count = test_opnum;
|
||||
run_ops = test_ops;
|
||||
run_opnames = test_opnames;
|
||||
}
|
||||
|
||||
@@ -9,15 +9,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) {
|
||||
@@ -85,7 +85,7 @@ testResult_t ReduceScatterRunTest(struct threadArgs* args, int root, ncclDataTyp
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
} else {
|
||||
type_count = ncclNumTypes;
|
||||
type_count = test_typenum;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
@@ -95,7 +95,7 @@ testResult_t ReduceScatterRunTest(struct threadArgs* args, int root, ncclDataTyp
|
||||
run_opnames = &opName;
|
||||
op_count = 1;
|
||||
} else {
|
||||
op_count = sizeof(test_ops)/sizeof(test_ops[0]);
|
||||
op_count = test_opnum;
|
||||
run_ops = test_ops;
|
||||
run_opnames = test_opnames;
|
||||
}
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Copyright (c) 2016-2021, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2020-2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
@@ -8,33 +8,29 @@
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "common.h"
|
||||
|
||||
//#define DEBUG_PRINT
|
||||
#define USE_RCCL_GATHER_SCATTER
|
||||
|
||||
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", "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 %6s", size, count, typeName, opName);
|
||||
PRINT("%12li %12li %8s %6i", size, count, typeName, root);
|
||||
}
|
||||
|
||||
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;
|
||||
*recvInplaceOffset = count/nranks;
|
||||
*paramcount = count/nranks;
|
||||
}
|
||||
|
||||
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; i<args->nGpus; i++) {
|
||||
int gpuid = args->localRank*args->nThreads*args->nGpus + args->thread*args->nGpus + i;
|
||||
@@ -42,31 +38,15 @@ testResult_t ScatterInitData(struct threadArgs* args, ncclDataType_t type, ncclR
|
||||
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; j<nranks; j++) {
|
||||
TESTCHECK(InitData(((char*)data)+args->expectedBytes*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; j<args->sendBytes/sizeof(int); j++) {
|
||||
printf("%d:%d ", j, dataHost[j]);
|
||||
}
|
||||
printf("\n");
|
||||
free(dataHost);
|
||||
#endif
|
||||
}
|
||||
TESTCHECK(InitData(args->expected[i], recvcount, type, rep, rank));
|
||||
if (rank == root) TESTCHECK(InitData(data, sendcount, type, rep, rank));
|
||||
TESTCHECK(InitData(args->expected[i], recvcount, type, rep+rank*recvcount, root));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
// We don't support in-place scatter
|
||||
args->reportErrors = in_place ? 0 : 1;
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void ScatterGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw, int nranks) {
|
||||
double baseBw = (double)(count * typesize) / 1.0E9 / sec;
|
||||
double baseBw = (double)(count * nranks * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double)(nranks-1))/((double)(nranks));
|
||||
@@ -76,31 +56,21 @@ void ScatterGetBw(size_t count, int typesize, double sec, double* algBw, double*
|
||||
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));
|
||||
int rank;
|
||||
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
||||
size_t rankOffset = count * wordSize(type);
|
||||
if (count == 0) return testSuccess;
|
||||
|
||||
int rank;
|
||||
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
||||
#if NCCL_MAJOR < 2 || NCCL_MINOR < 7
|
||||
printf("NCCL 2.7 or later is needed for scatter. This test was compiled with %d.%d.\n", NCCL_MAJOR, NCCL_MINOR);
|
||||
return testNcclError;
|
||||
#else
|
||||
#if defined(RCCL_GATHER_SCATTER) && defined(USE_RCCL_GATHER_SCATTER)
|
||||
if (rank == root)
|
||||
NCCLCHECK(ncclScatter(sendbuff, recvbuff, count, type, root, comm, stream));
|
||||
else
|
||||
NCCLCHECK(ncclScatter(0, recvbuff, count, type, root, comm, stream));
|
||||
#else
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
if (rank == root) {
|
||||
for (int r=0; r<nRanks; r++)
|
||||
for (int r=0; r<nRanks; r++) {
|
||||
NCCLCHECK(ncclSend(((char*)sendbuff)+r*rankOffset, count, type, r, comm, stream));
|
||||
}
|
||||
}
|
||||
NCCLCHECK(ncclRecv(recvbuff, count, type, root, comm, stream));
|
||||
NCCLCHECK(ncclGroupEnd());
|
||||
#endif
|
||||
|
||||
return testSuccess;
|
||||
#endif
|
||||
}
|
||||
|
||||
struct testColl scatterTest = {
|
||||
@@ -121,19 +91,29 @@ testResult_t ScatterRunTest(struct threadArgs* args, int root, ncclDataType_t ty
|
||||
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 = ncclNumTypes;
|
||||
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<type_count; i++) {
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t)0, "", root));
|
||||
for (int j=begin_root; j<=end_root; j++) {
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t)0, "", j));
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Copyright (c) 2016-2021, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2020-2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
@@ -9,15 +9,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 SendRecvGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) {
|
||||
@@ -37,12 +37,15 @@ testResult_t SendRecvInitData(struct threadArgs* args, ncclDataType_t type, nccl
|
||||
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);
|
||||
TESTCHECK(InitData(args->sendbuffs[i], sendcount, type, rep, rank));
|
||||
TESTCHECK(InitData(args->recvbuffs[i], recvcount, type, rep, rank));
|
||||
int src = rank < nranks/2 ? rank : rank - nranks/2;
|
||||
TESTCHECK(InitData(args->expected[i], recvcount, type, rep, src));
|
||||
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));
|
||||
int peer = (rank-1+nranks)%nranks;
|
||||
TESTCHECK(InitData(args->expected[i], recvcount, type, rep, peer));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
// We don't support in-place sendrecv
|
||||
args->reportErrors = in_place ? 0 : 1;
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
@@ -50,28 +53,26 @@ void SendRecvGetBw(size_t count, int typesize, double sec, double* algBw, double
|
||||
double baseBw = (double)(count * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = nranks/2;
|
||||
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, hipStream_t stream) {
|
||||
int rank, peer, nranks, npairs;
|
||||
int nRanks;
|
||||
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
||||
int rank;
|
||||
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
||||
NCCLCHECK(ncclCommCount(comm, &nranks));
|
||||
npairs = nranks / 2;
|
||||
#if NCCL_MAJOR >= 2 && NCCL_MINOR >= 7
|
||||
if (rank < npairs) {
|
||||
peer = rank + npairs;
|
||||
NCCLCHECK(ncclSend(sendbuff, count, type, peer, comm, stream));
|
||||
} else if (rank < 2*npairs) {
|
||||
peer = rank - npairs;
|
||||
NCCLCHECK(ncclRecv(recvbuff, count, type, peer, comm, stream));
|
||||
}
|
||||
#endif
|
||||
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 = {
|
||||
struct testColl sendRecvTest = {
|
||||
"SendRecv",
|
||||
SendRecvGetCollByteCount,
|
||||
SendRecvInitData,
|
||||
@@ -85,23 +86,36 @@ void SendRecvGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int
|
||||
}
|
||||
|
||||
testResult_t SendRecvRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
args->collTest = &sendrecvTest;
|
||||
args->collTest = &sendRecvTest;
|
||||
ncclDataType_t *run_types;
|
||||
const char **run_typenames;
|
||||
int type_count;
|
||||
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 = ncclNumTypes;
|
||||
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<type_count; i++) {
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t)0, "", 0));
|
||||
for (int j=0; j<op_count; j++) {
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], run_ops[j], run_opnames[j], -1));
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
In neuem Issue referenzieren
Einen Benutzer sperren