Merge pull request #116 from nileshnegi/sync/nccl-tests/02-28-2025
[SYNC] NCCL-Tests v2.14.1
[ROCm/rccl-tests commit: 83d38d91b6]
이 커밋은 다음에 포함됨:
Nilesh M Negi
GitHub
@@ -128,7 +128,7 @@ add_library(rccl_common OBJECT ${HIP_COMMON_SOURCES})
|
||||
add_dependencies(rccl_common hipify git_version_check)
|
||||
target_link_libraries(rccl_common roc::rccl hip::device)
|
||||
if(USE_MPI)
|
||||
target_link_libraries(rccl_common MPI::MPI_CXX)
|
||||
target_link_libraries(rccl_common MPI::MPI_CXX)
|
||||
endif()
|
||||
|
||||
# Compile tests
|
||||
|
||||
@@ -1,32 +1,96 @@
|
||||
#
|
||||
# Copyright (c) 2015-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
# Modifications are Copyright (c) 2019-2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
# Modifications are Copyright (c) 2019-2025 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# See LICENSE.txt for license information
|
||||
#
|
||||
|
||||
ROCM_PATH ?= /opt/rocm
|
||||
MPI_HOME ?= /usr/lib/openmpi
|
||||
MPI_HOME ?= /usr/lib/x86_64-linux-gnu
|
||||
PREFIX ?= /usr/local
|
||||
VERBOSE ?= 0
|
||||
DEBUG ?= 0
|
||||
NCCL_HOME ?= ""
|
||||
CUSTOM_RCCL_LIB ?= ""
|
||||
|
||||
HIPCC = $(ROCM_PATH)/bin/hipcc
|
||||
HIPCC = $(ROCM_PATH)/bin/amdclang++
|
||||
HIPCONFIG = $(ROCM_PATH)/bin/hipconfig
|
||||
CXX = $(HIPCC)
|
||||
|
||||
HIPCUFLAGS := -std=c++14
|
||||
LDFLAGS :=
|
||||
HIPLDFLAGS :=
|
||||
|
||||
HIP_VERSION = $(strip $(shell which $(HIPCONFIG) >/dev/null && $(HIPCONFIG) --version))
|
||||
HIP_MAJOR = $(shell echo $(HIP_VERSION) | cut -d "." -f 1)
|
||||
HIP_MINOR = $(shell echo $(HIP_VERSION) | cut -d "." -f 2)
|
||||
|
||||
# Better define GPU_TARGETS in your environment to the minimal set
|
||||
# of archs to reduce compile time.
|
||||
# Currently, supports gfx906,gfx908,gfx90a,gfx942,gfx950,gfx1030,gfx1100,gfx1101,gfx1102,gfx1200,gfx1201
|
||||
ifndef GPU_TARGETS
|
||||
GPU_TARGETS = gfx906 gfx908 gfx90a
|
||||
ifeq ($(shell test "0$(HIP_MAJOR)" -eq 6; echo $$?),0)
|
||||
# Include gfx942 support if we're using ROCm 6.0 or above
|
||||
GPU_TARGETS += gfx942
|
||||
ifeq ($(shell test "0$(HIP_MINOR)" -ge 5; echo $$?),0)
|
||||
# Include gfx950 support if we're using ROCm 6.5 or above
|
||||
GPU_TARGETS += gfx950
|
||||
endif
|
||||
endif
|
||||
GPU_TARGETS += gfx1030 gfx1100 gfx1101 gfx1102 gfx1200 gfx1201
|
||||
endif
|
||||
|
||||
GPU_TARGETS_FLAGS = $(foreach target,$(GPU_TARGETS),"--offload-arch=$(target)")
|
||||
|
||||
#CUDA_VERSION = $(strip $(shell which $(NVCC) >/dev/null && $(NVCC) --version | grep release | sed 's/.*release //' | sed 's/\,.*//'))
|
||||
#CUDA_MAJOR = $(shell echo $(CUDA_VERSION) | cut -d "." -f 1)
|
||||
#CUDA_MINOR = $(shell echo $(CUDA_VERSION) | cut -d "." -f 2)
|
||||
#
|
||||
## Better define NVCC_GENCODE in your environment to the minimal set
|
||||
## of archs to reduce compile time.
|
||||
#ifeq ($(shell test "0$(CUDA_MAJOR)" -eq 12 -a "0$(CUDA_MINOR)" -ge 8 -o "0$(CUDA_MAJOR)" -ge 13; echo $$?),0)
|
||||
## Include Blackwell support if we're using CUDA12.8 or above
|
||||
#NVCC_GENCODE ?= -gencode=arch=compute_80,code=sm_80 \
|
||||
# -gencode=arch=compute_90,code=sm_90 \
|
||||
# -gencode=arch=compute_100,code=sm_100 \
|
||||
# -gencode=arch=compute_120,code=sm_120 \
|
||||
# -gencode=arch=compute_120,code=compute_120
|
||||
#else ifeq ($(shell test "0$(CUDA_MAJOR)" -ge 12; echo $$?),0)
|
||||
#NVCC_GENCODE ?= -gencode=arch=compute_60,code=sm_60 \
|
||||
# -gencode=arch=compute_61,code=sm_61 \
|
||||
# -gencode=arch=compute_70,code=sm_70 \
|
||||
# -gencode=arch=compute_80,code=sm_80 \
|
||||
# -gencode=arch=compute_90,code=sm_90 \
|
||||
# -gencode=arch=compute_90,code=compute_90
|
||||
#else ifeq ($(shell test "0$(CUDA_MAJOR)" -ge 11; echo $$?),0)
|
||||
#NVCC_GENCODE ?= -gencode=arch=compute_60,code=sm_60 \
|
||||
# -gencode=arch=compute_61,code=sm_61 \
|
||||
# -gencode=arch=compute_70,code=sm_70 \
|
||||
# -gencode=arch=compute_80,code=sm_80 \
|
||||
# -gencode=arch=compute_80,code=compute_80
|
||||
#else
|
||||
#NVCC_GENCODE ?= -gencode=arch=compute_35,code=sm_35 \
|
||||
# -gencode=arch=compute_50,code=sm_50 \
|
||||
# -gencode=arch=compute_60,code=sm_60 \
|
||||
# -gencode=arch=compute_61,code=sm_61 \
|
||||
# -gencode=arch=compute_70,code=sm_70 \
|
||||
# -gencode=arch=compute_70,code=compute_70
|
||||
#endif
|
||||
|
||||
ifneq ($(NCCL_HOME), "")
|
||||
HIPCUFLAGS += -I$(NCCL_HOME)/ -I$(NCCL_HOME)/include
|
||||
HIPLDFLAGS += -Wl,-rpath,$(NCCL_HOME) -L$(NCCL_HOME) -L$(NCCL_HOME)/lib
|
||||
HIPLDFLAGS += -Wl,-rpath,$(NCCL_HOME) -L$(NCCL_HOME) -L$(NCCL_HOME)/lib
|
||||
endif
|
||||
|
||||
HIPCUFLAGS += -I$(ROCM_PATH)/include
|
||||
HIPCUFLAGS += -I$(ROCM_PATH)/include/hip
|
||||
HIPCUFLAGS += -x hip -D__HIP_PLATFORM_AMD__ -D__HIPCC__ $(GPU_TARGETS_FLAGS)
|
||||
LDFLAGS += -L$(ROCM_PATH)/lib -lhsa-runtime64 -lrt
|
||||
HIPLDFLAGS += -L$(CUSTOM_RCCL_LIB) -L$(ROCM_PATH)/lib -lhsa-runtime64 -lrt -pthread
|
||||
ifneq ($(CUSTOM_RCCL_LIB), "")
|
||||
HIPLDFLAGS += -L$(CUSTOM_RCCL_LIB)
|
||||
endif
|
||||
HIPLDFLAGS += -L$(ROCM_PATH)/lib -lhsa-runtime64 -lamdhip64 -lstdc++ -lrt -pthread
|
||||
|
||||
ifeq ($(DEBUG), 0)
|
||||
HIPCUFLAGS += -O3
|
||||
@@ -46,20 +110,20 @@ HIPIFY_DIR ?= $(BUILDDIR)/hipify
|
||||
.PRECIOUS: $(HIPIFY_DIR)/%.cu.cpp $(HIPIFY_DIR)/%.h
|
||||
|
||||
ifeq ($(MPI), 1)
|
||||
HIPCUFLAGS += -DMPI_SUPPORT -I${MPI_HOME}/include -I${MPI_HOME}/include/mpi
|
||||
HIPLDFLAGS += -L${MPI_HOME}/lib -lmpi
|
||||
HIPCUFLAGS += -DMPI_SUPPORT -I${MPI_HOME}/include -I${MPI_HOME}/include/openmpi -I${MPI_HOME}/openmpi/include -I${MPI_HOME}/openmpi/include/openmpi
|
||||
HIPLDFLAGS += -L${MPI_HOME}/lib -L${MPI_HOME}/openmpi/lib -lmpi
|
||||
else ifeq ($(MPICH), 1)
|
||||
HIPCUFLAGS += -DMPI_SUPPORT -I/usr/include/mpich -I/usr/include/x86_64-linux-gnu/mpich
|
||||
HIPLDFLAGS += -L/usr/lib -lmpich
|
||||
HIPCUFLAGS += -DMPI_SUPPORT -I${MPI_HOME}/include -I${MPI_HOME}/mpich/include -I/usr/include/x86_64-linux-gnu/mpich
|
||||
HIPLDFLAGS += -L${MPI_HOME}/lib -L${MPI_HOME}/mpich/lib -lmpich
|
||||
endif
|
||||
|
||||
LIBRARIES += rccl
|
||||
HIPLDFLAGS += $(LIBRARIES:%=-l%)
|
||||
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 scatter gather sendrecv alltoallv
|
||||
BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall scatter gather sendrecv alltoallv hypercube
|
||||
BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf)
|
||||
|
||||
GIT_VERSION_FILE := ${DST_DIR}/src/git_version.cpp
|
||||
@@ -94,8 +158,8 @@ ${HIPIFY_DIR}/%.h: %.h
|
||||
${DST_DIR}/%.o: ${HIPIFY_DIR}/%.cu.cpp ${HIPIFY_DIR}/common.h $(TEST_VERIFIABLE_HDRS) $(GIT_VERSION_FILE)
|
||||
@printf "Compiling %-35s > %s\n" $< $@
|
||||
@mkdir -p ${DST_DIR}
|
||||
echo "$(HIPCC) -o $@ $(HIPCUFLAGS) -I. -c $<"
|
||||
$(HIPCC) -o $@ $(HIPCUFLAGS) -I. -c $<
|
||||
echo "$(HIPCC) $(HIPCUFLAGS) -I. -c -o $@ $<"
|
||||
$(HIPCC) $(HIPCUFLAGS) -I. -c -o $@ $<
|
||||
|
||||
${DST_DIR}/timer.o: timer.cc timer.h
|
||||
@printf "Compiling %-35s > %s\n" $< $@
|
||||
@@ -105,6 +169,6 @@ ${DST_DIR}/timer.o: timer.cc timer.h
|
||||
${DST_DIR}/%_perf:${DST_DIR}/%.o ${DST_DIR}/common.o ${DST_DIR}/timer.o $(TEST_VERIFIABLE_OBJS) $(DST_DIR)/src/git_version.cpp
|
||||
@printf "Linking %-35s > %s\n" $< $@
|
||||
@mkdir -p ${DST_DIR}
|
||||
echo "$(HIPCC) -o $@ $(HIPCUFLAGS) $^ ${HIPLDFLAGS}"
|
||||
$(HIPCC) -o $@ $(HIPCUFLAGS) $^ ${HIPLDFLAGS}
|
||||
echo "$(HIPCC) -o $@ $^ $(HIPLDFLAGS)"
|
||||
$(HIPCC) -o $@ $^ $(HIPLDFLAGS)
|
||||
|
||||
|
||||
@@ -8,10 +8,8 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
#define ALIGN 4
|
||||
|
||||
void AllGatherGetCollByteCount(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;
|
||||
void AllGatherGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) {
|
||||
size_t base = (count/nranks) & -(16/eltSize);
|
||||
*sendcount = base;
|
||||
*recvcount = base*nranks;
|
||||
*sendInplaceOffset = base;
|
||||
@@ -61,7 +59,7 @@ struct testColl allGatherTest = {
|
||||
|
||||
void AllGatherGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
AllGatherGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
AllGatherGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t AllGatherRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -8,7 +8,7 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
void AllReduceGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) {
|
||||
void AllReduceGetCollByteCount(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;
|
||||
*recvcount = count;
|
||||
*sendInplaceOffset = 0;
|
||||
@@ -56,7 +56,7 @@ struct testColl allReduceTest = {
|
||||
|
||||
void AllReduceGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
AllReduceGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
AllReduceGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t AllReduceRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -8,12 +8,12 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
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;
|
||||
void AlltoAllGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) {
|
||||
*paramcount = (count/nranks) & -(16/eltSize);
|
||||
*sendcount = nranks*(*paramcount);
|
||||
*recvcount = *sendcount;
|
||||
*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) {
|
||||
@@ -61,7 +61,7 @@ struct testColl alltoAllTest = {
|
||||
|
||||
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);
|
||||
AlltoAllGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t AlltoAllRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -10,7 +10,7 @@
|
||||
|
||||
#define USE_RCCL_GATHER_SCATTER
|
||||
|
||||
void AlltoAllvGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) {
|
||||
void AlltoAllvGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) {
|
||||
if (count < nranks*nranks/2) {
|
||||
*sendcount = 0;
|
||||
*recvcount = 0;
|
||||
@@ -18,11 +18,11 @@ void AlltoAllvGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *par
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = 0;
|
||||
} else {
|
||||
*sendcount = (count/nranks)*nranks;
|
||||
*recvcount = (count/nranks)*nranks;
|
||||
*paramcount = (count/nranks) & -(16/eltSize);
|
||||
*sendcount = nranks*(*paramcount);
|
||||
*recvcount = *sendcount;
|
||||
*sendInplaceOffset = 0;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = count/nranks;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -161,7 +161,7 @@ struct testColl alltoAllTest = {
|
||||
|
||||
void AlltoAllvGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
AlltoAllvGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
AlltoAllvGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t AlltoAllvRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -8,7 +8,7 @@
|
||||
#include "cuda_runtime.h"
|
||||
#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) {
|
||||
void BroadcastGetCollByteCount(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;
|
||||
*recvcount = count;
|
||||
*sendInplaceOffset = 0;
|
||||
@@ -65,7 +65,7 @@ struct testColl broadcastTest = {
|
||||
|
||||
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);
|
||||
BroadcastGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t BroadcastRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
+179
-117
@@ -16,6 +16,8 @@
|
||||
#include <type_traits>
|
||||
#include <getopt.h>
|
||||
#include <libgen.h>
|
||||
#include <string.h>
|
||||
#include <ctype.h>
|
||||
#include "cuda.h"
|
||||
#include <vector>
|
||||
#include <utility>
|
||||
@@ -90,13 +92,13 @@ static int datacheck = 1;
|
||||
static int warmup_iters = 5;
|
||||
static int iters = 20;
|
||||
static int agg_iters = 1;
|
||||
static int run_cycles = 1;
|
||||
static int ncclop = ncclSum;
|
||||
static int nccltype = ncclFloat;
|
||||
static int ncclroot = 0;
|
||||
static int parallel_init = 0;
|
||||
static int blocking_coll = 0;
|
||||
static int memorytype = 0;
|
||||
static int stress_cycles = 1;
|
||||
static uint32_t cumask[4];
|
||||
static int streamnull = 0;
|
||||
static int timeout = 0;
|
||||
@@ -121,6 +123,7 @@ Reporter::Reporter(std::string fileName, std::string outputFormat) : _outputForm
|
||||
_out = std::ofstream(fileName, std::ios_base::out);
|
||||
_outputValid = true;
|
||||
if (_outputFormat == "csv") {
|
||||
_out << "numCycle, ";
|
||||
_out << "collective, ";
|
||||
#ifdef MPI_SUPPORT
|
||||
_out << "ranks, rankspernode, gpusperrank, ";
|
||||
@@ -133,10 +136,11 @@ Reporter::Reporter(std::string fileName, std::string outputFormat) : _outputForm
|
||||
}
|
||||
}
|
||||
|
||||
void Reporter::setParameters(const char* name, const char* typeName, const char* opName) {
|
||||
void Reporter::setParameters(const size_t numCycle, const char* name, const char* typeName, const char* opName) {
|
||||
if (!isMainThread() || !_outputValid)
|
||||
return;
|
||||
|
||||
_numCycle = numCycle;
|
||||
_collectiveName = name;
|
||||
_typeName = typeName;
|
||||
_opName = opName;
|
||||
@@ -150,6 +154,7 @@ void Reporter::addResult(int gpusPerRank, int ranksPerNode, int totalRanks, size
|
||||
std::string wrongEltsStr = (wrongElts == -1) ? "N/A" : std::to_string(wrongElts);
|
||||
int nodes = totalRanks / ranksPerNode;
|
||||
|
||||
outputValuesKeys.push_back(makeValueKeyPair(_numCycle, "numCycle"));
|
||||
outputValuesKeys.push_back(makeValueKeyPair(_collectiveName, "name"));
|
||||
#ifdef MPI_SUPPORT
|
||||
outputValuesKeys.push_back(makeValueKeyPair(nodes, "nodes"));
|
||||
@@ -614,8 +619,8 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
|
||||
Barrier(args);
|
||||
|
||||
#if HIP_VERSION >= 50221310
|
||||
cudaGraph_t graphs[args->nGpus];
|
||||
cudaGraphExec_t graphExec[args->nGpus];
|
||||
std::vector<cudaGraph_t> graphs(args->nGpus);
|
||||
std::vector<cudaGraphExec_t> graphExec(args->nGpus);
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
// Begin cuda graph capture
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
@@ -642,11 +647,11 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
// End cuda graph capture
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i));
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs.data()+i));
|
||||
}
|
||||
// Instantiate cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0));
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec.data()+i, graphs[i], NULL, NULL, 0));
|
||||
}
|
||||
// Resync CPU, restart timing, launch cuda graph
|
||||
Barrier(args);
|
||||
@@ -705,11 +710,11 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
// End cuda graph capture
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i));
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs.data()+i));
|
||||
}
|
||||
// Instantiate cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0));
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec.data()+i, graphs[i], NULL, NULL, 0));
|
||||
}
|
||||
// Launch cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
@@ -774,7 +779,7 @@ void setupArgs(size_t size, ncclDataType_t type, struct threadArgs* args) {
|
||||
size_t count, sendCount, recvCount, paramCount, sendInplaceOffset, recvInplaceOffset;
|
||||
|
||||
count = size / wordSize(type);
|
||||
args->collTest->getCollByteCount(&sendCount, &recvCount, ¶mCount, &sendInplaceOffset, &recvInplaceOffset, (size_t)count, (size_t)nranks);
|
||||
args->collTest->getCollByteCount(&sendCount, &recvCount, ¶mCount, &sendInplaceOffset, &recvInplaceOffset, (size_t)count, wordSize(type), (size_t)nranks);
|
||||
|
||||
args->nbytes = paramCount * wordSize(type);
|
||||
args->sendBytes = sendCount * wordSize(type);
|
||||
@@ -790,8 +795,8 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
|
||||
// Warm-up for large size
|
||||
setupArgs(args->maxbytes, type, args);
|
||||
#if HIP_VERSION >= 50221310
|
||||
cudaGraph_t graphs[args->nGpus];
|
||||
cudaGraphExec_t graphExec[args->nGpus];
|
||||
std::vector<cudaGraph_t> graphs(args->nGpus);
|
||||
std::vector<cudaGraphExec_t> graphExec(args->nGpus);
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
// Begin cuda graph capture
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
@@ -811,11 +816,11 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
// End cuda graph capture
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i));
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs.data()+i));
|
||||
}
|
||||
// Instantiate cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0));
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec.data()+i, graphs[i], NULL, NULL, 0));
|
||||
}
|
||||
// Resync CPU, restart timing, launch cuda graph
|
||||
Barrier(args);
|
||||
@@ -861,11 +866,11 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
|
||||
if (cudaGraphLaunches >= 1) {
|
||||
// End cuda graph capture
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i));
|
||||
CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs.data()+i));
|
||||
}
|
||||
// Instantiate cuda graph
|
||||
for (int i=0; i<args->nGpus; i++) {
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0));
|
||||
CUDACHECK(cudaGraphInstantiate(graphExec.data()+i, graphs[i], NULL, NULL, 0));
|
||||
}
|
||||
// Resync CPU, restart timing, launch cuda graph
|
||||
Barrier(args);
|
||||
@@ -889,26 +894,31 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
|
||||
}
|
||||
#endif
|
||||
|
||||
if (args->reporter) {
|
||||
args->reporter->setParameters(args->collTest->name, typeName, opName);
|
||||
}
|
||||
// Benchmark
|
||||
long repeat = run_cycles;
|
||||
size_t iter = 0;
|
||||
|
||||
for (size_t iter = 0; iter < stress_cycles; iter++) {
|
||||
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);
|
||||
char rootName[100];
|
||||
sprintf(rootName, "%6i", root);
|
||||
PRINT("%12li %12li %8s %6s %6s", std::max(args->sendBytes, args->expectedBytes), args->nbytes / wordSize(type), typeName, opName, rootName);
|
||||
if (enable_out_of_place) {
|
||||
TESTCHECK(BenchTime(args, type, op, root, 0));
|
||||
usleep(delay_inout_place);
|
||||
}
|
||||
TESTCHECK(BenchTime(args, type, op, root, 1));
|
||||
PRINT("\n");
|
||||
do {
|
||||
if (run_cycles > 1) PRINT("# Testing %lu cycle.\n", iter+1);
|
||||
if (args->reporter) {
|
||||
args->reporter->setParameters(iter, args->collTest->name, typeName, opName);
|
||||
}
|
||||
}
|
||||
for (size_t size = args->minbytes; size<=args->maxbytes; size = ((args->stepfactor > 1) ? size*args->stepfactor : size+args->stepbytes)) {
|
||||
setupArgs(size, type, args);
|
||||
char rootName[100];
|
||||
sprintf(rootName, "%6i", root);
|
||||
PRINT("%12li %12li %8s %6s %6s", std::max(args->sendBytes, args->expectedBytes), args->nbytes / wordSize(type), typeName, opName, rootName);
|
||||
if (enable_out_of_place) {
|
||||
TESTCHECK(BenchTime(args, type, op, root, 0));
|
||||
usleep(delay_inout_place);
|
||||
}
|
||||
TESTCHECK(BenchTime(args, type, op, root, 1));
|
||||
PRINT("\n");
|
||||
}
|
||||
--repeat;
|
||||
++iter;
|
||||
} while(repeat != 0);
|
||||
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
@@ -1052,26 +1062,27 @@ int main(int argc, char* argv[]) {
|
||||
{"iters", required_argument, 0, 'n'},
|
||||
{"agg_iters", required_argument, 0, 'm'},
|
||||
{"warmup_iters", required_argument, 0, 'w'},
|
||||
{"run_cycles", required_argument, 0, 'N'},
|
||||
{"parallel_init", required_argument, 0, 'p'},
|
||||
{"check", required_argument, 0, 'c'},
|
||||
{"op", required_argument, 0, 'o'},
|
||||
{"datatype", required_argument, 0, 'd'},
|
||||
{"root", required_argument, 0, 'r'},
|
||||
{"blocking", required_argument, 0, 'z'},
|
||||
{"memory_type", required_argument, 0, 'y'}, //RCCL
|
||||
{"stress_cycles", required_argument, 0, 's'}, //RCCL
|
||||
{"cumask", required_argument, 0, 'u'}, //RCCL
|
||||
{"stream_null", required_argument, 0, 'y'}, //NCCL
|
||||
{"timeout", required_argument, 0, 'T'}, //NCCL
|
||||
{"stream_null", required_argument, 0, 'y'},
|
||||
{"timeout", required_argument, 0, 'T'},
|
||||
{"cudagraph", required_argument, 0, 'G'},
|
||||
{"report_cputime", required_argument, 0, 'C'},
|
||||
{"average", required_argument, 0, 'a'},
|
||||
{"out_of_place", required_argument, 0, 'O'},
|
||||
{"cache_flush", required_argument, 0, 'F'},
|
||||
{"rotating_tensor", required_argument, 0, 'E'},
|
||||
{"local_register", required_argument, 0, 'R'},
|
||||
{"output_file", required_argument, 0, 'x'},
|
||||
{"output_format", required_argument, 0, 'Z'},
|
||||
{"memory_type", required_argument, 0, 'y'}, //RCCL
|
||||
{"cumask", required_argument, 0, 'u'}, //RCCL
|
||||
{"out_of_place", required_argument, 0, 'O'}, //RCCL
|
||||
{"delay_inout_place", required_argument, 0, 'q'}, //RCCL
|
||||
{"cache_flush", required_argument, 0, 'F'}, //RCCL
|
||||
{"rotating_tensor", required_argument, 0, 'E'}, //RCCL
|
||||
{"output_file", required_argument, 0, 'x'}, //RCCL
|
||||
{"output_format", required_argument, 0, 'Z'}, //RCCL
|
||||
{"help", no_argument, 0, 'h'},
|
||||
{}
|
||||
};
|
||||
@@ -1079,7 +1090,7 @@ int main(int argc, char* argv[]) {
|
||||
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:T:G:C:O:F:E:R:a:y:s:u:h:q:x:Z:", longopts, &longindex);
|
||||
c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:N:p:c:o:d:r:z:y:T:G:C:a:R:Y:u:O:q:F:E:x:Z:h", longopts, &longindex);
|
||||
|
||||
if (c == -1)
|
||||
break;
|
||||
@@ -1108,7 +1119,12 @@ int main(int argc, char* argv[]) {
|
||||
maxBytes = (size_t)parsed;
|
||||
break;
|
||||
case 'i':
|
||||
stepBytes = strtol(optarg, NULL, 0);
|
||||
parsed = parsesize(optarg);
|
||||
if (parsed < 0) {
|
||||
fprintf(stderr, "invalid size specified for 'stepBytes'\n");
|
||||
return -1;
|
||||
}
|
||||
stepBytes = (size_t)parsed;
|
||||
break;
|
||||
case 'f':
|
||||
stepFactor = strtol(optarg, NULL, 0);
|
||||
@@ -1126,12 +1142,15 @@ int main(int argc, char* argv[]) {
|
||||
case 'w':
|
||||
warmup_iters = (int)strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'c':
|
||||
datacheck = (int)strtol(optarg, NULL, 0);
|
||||
case 'N':
|
||||
run_cycles = (int)strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'p':
|
||||
parallel_init = (int)strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'c':
|
||||
datacheck = (int)strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'o':
|
||||
ncclop = ncclstringtoop(optarg);
|
||||
break;
|
||||
@@ -1144,22 +1163,6 @@ int main(int argc, char* argv[]) {
|
||||
case 'z':
|
||||
blocking_coll = strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'Y':
|
||||
memorytype = ncclstringtomtype(optarg);
|
||||
break;
|
||||
case 's':
|
||||
stress_cycles = strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'u':
|
||||
{
|
||||
int nmasks = 0;
|
||||
char *mask = strtok(optarg, ",");
|
||||
while (mask != NULL && nmasks < 4) {
|
||||
cumask[nmasks++] = strtol(mask, NULL, 16);
|
||||
mask = strtok(NULL, ",");
|
||||
};
|
||||
}
|
||||
break;
|
||||
case 'y':
|
||||
streamnull = strtol(optarg, NULL, 0);
|
||||
break;
|
||||
@@ -1176,9 +1179,37 @@ int main(int argc, char* argv[]) {
|
||||
case 'C':
|
||||
report_cputime = strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'a':
|
||||
average = (int)strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'R':
|
||||
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
|
||||
if ((int)strtol(optarg, NULL, 0)) {
|
||||
local_register = 1;
|
||||
}
|
||||
#else
|
||||
printf("Option -R (register) is not supported before NCCL 2.19. Ignoring\n");
|
||||
#endif
|
||||
break;
|
||||
case 'Y':
|
||||
memorytype = ncclstringtomtype(optarg);
|
||||
break;
|
||||
case 'u':
|
||||
{
|
||||
int nmasks = 0;
|
||||
char *mask = strtok(optarg, ",");
|
||||
while (mask != NULL && nmasks < 4) {
|
||||
cumask[nmasks++] = strtol(mask, NULL, 16);
|
||||
mask = strtok(NULL, ",");
|
||||
};
|
||||
}
|
||||
break;
|
||||
case 'O':
|
||||
enable_out_of_place = strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'q':
|
||||
delay_inout_place = (int)strtol(optarg, NULL, 10);
|
||||
break;
|
||||
case 'F':
|
||||
enable_cache_flush = strtol(optarg, NULL, 0);
|
||||
if (enable_cache_flush > 0) {
|
||||
@@ -1190,20 +1221,6 @@ int main(int argc, char* argv[]) {
|
||||
case 'E':
|
||||
enable_rotating_tensor = strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'a':
|
||||
average = (int)strtol(optarg, NULL, 0);
|
||||
break;
|
||||
case 'q':
|
||||
delay_inout_place = (int)strtol(optarg, NULL, 10);
|
||||
case 'R':
|
||||
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
|
||||
if ((int)strtol(optarg, NULL, 0)) {
|
||||
local_register = 1;
|
||||
}
|
||||
#else
|
||||
printf("Option -R (register) is not supported before NCCL 2.19. Ignoring\n");
|
||||
#endif
|
||||
break;
|
||||
case 'x':
|
||||
output_file = optarg;
|
||||
break;
|
||||
@@ -1223,6 +1240,7 @@ int main(int argc, char* argv[]) {
|
||||
"[-n,--iters <iteration count>] \n\t"
|
||||
"[-m,--agg_iters <aggregated iteration count>] \n\t"
|
||||
"[-w,--warmup_iters <warmup iteration count>] \n\t"
|
||||
"[-N,--run_cycles <cycle count> run & print each cycle (default: 1; 0=infinite)] \n\t"
|
||||
"[-p,--parallel_init <0/1>] \n\t"
|
||||
"[-c,--check <check iteration count>] \n\t"
|
||||
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0)
|
||||
@@ -1235,19 +1253,18 @@ int main(int argc, char* argv[]) {
|
||||
"[-d,--datatype <nccltype/all>] \n\t"
|
||||
"[-r,--root <root/all>] \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"
|
||||
"[-y,--stream_null <0/1>] \n\t"
|
||||
"[-T,--timeout <time in seconds>] \n\t"
|
||||
"[-G,--cudagraph <num graph launches>] \n\t"
|
||||
"[-C,--report_cputime <0/1>] \n\t"
|
||||
"[-a,--average <0/1/2/3> report average iteration time <0=RANK0/1=AVG/2=MIN/3=MAX>] \n\t"
|
||||
"[-R,--local_register <1/0> enable local buffer registration on send/recv buffers (default: disable)] \n\t"
|
||||
"[-Y,--memory_type <coarse/fine/host/managed>] \n\t"
|
||||
"[-u,--cumask <d0,d1,d2,d3>] \n\t"
|
||||
"[-O,--out_of_place <0/1>] \n\t"
|
||||
"[-q,--delay <delay between out-of-place and in-place in microseconds>] \n\t"
|
||||
"[-F,--cache_flush <number of iterations between instruction cache flush>] \n\t"
|
||||
"[-E,--rotating_tensor <0/1>] \n\t"
|
||||
"[-a,--average <0/1/2/3> report average iteration time <0=RANK0/1=AVG/2=MIN/3=MAX>] \n\t"
|
||||
"[-q,--delay <delay between out-of-place and in-place in microseconds>] \n\t"
|
||||
"[-R,--local_register <1/0> enable local buffer registration on send/recv buffers (default: disable)] \n\t"
|
||||
"[-x,--output_file <output file name>] \n\t"
|
||||
"[-Z,--output_format <output format <csv|json>] \n\t"
|
||||
"[-h,--help]\n",
|
||||
@@ -1283,6 +1300,26 @@ int main(int argc, char* argv[]) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
#ifdef MPI_SUPPORT
|
||||
// parse int for base 2/10/16, will ignore first whitespaces
|
||||
static bool parseInt(char *s, int *num) {
|
||||
char *p = NULL;
|
||||
if (!s || !num)
|
||||
return false;
|
||||
while (*s && isspace(*s)) ++s;
|
||||
if (!*s) return false;
|
||||
|
||||
if (strncasecmp(s, "0b", 2) == 0)
|
||||
*num = (int)strtoul(s + 2, &p, 2);
|
||||
else
|
||||
*num = (int)strtoul(s, &p, 0);
|
||||
|
||||
if (p == s)
|
||||
return false;
|
||||
return true;
|
||||
}
|
||||
#endif
|
||||
|
||||
testResult_t run() {
|
||||
int totalProcs = 1, proc = 0, ncclProcs = 1, ncclProc = 0, color = 0;
|
||||
int localRank = 0;
|
||||
@@ -1293,18 +1330,41 @@ testResult_t run() {
|
||||
#ifdef MPI_SUPPORT
|
||||
MPI_Comm_size(MPI_COMM_WORLD, &totalProcs);
|
||||
MPI_Comm_rank(MPI_COMM_WORLD, &proc);
|
||||
uint64_t hostHashs[totalProcs];
|
||||
std::vector<uint64_t> hostHashs(totalProcs);
|
||||
hostHashs[proc] = getHostHash(hostname);
|
||||
MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD);
|
||||
MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs.data(), sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD);
|
||||
for (int p=0; p<totalProcs; p++) {
|
||||
if (p == proc) break;
|
||||
if (hostHashs[p] == hostHashs[proc]) localRank++;
|
||||
}
|
||||
|
||||
char* str = getenv("NCCL_TESTS_SPLIT_MASK");
|
||||
uint64_t mask = str ? strtoul(str, NULL, 16) : 0;
|
||||
char *splitMaskEnv = NULL;
|
||||
if ((splitMaskEnv = getenv("NCCL_TESTS_SPLIT_MASK"))) {
|
||||
color = proc & strtoul(splitMaskEnv, NULL, 16);
|
||||
} else if ((splitMaskEnv = getenv("NCCL_TESTS_SPLIT"))) {
|
||||
if (
|
||||
(strncasecmp(splitMaskEnv, "AND", strlen("AND")) == 0 && parseInt(splitMaskEnv + strlen("AND"), &color)) ||
|
||||
(strncasecmp(splitMaskEnv, "&", strlen("&")) == 0 && parseInt(splitMaskEnv + strlen("&"), &color))
|
||||
)
|
||||
color = proc & color;
|
||||
if (
|
||||
(strncasecmp(splitMaskEnv, "OR", strlen("OR")) == 0 && parseInt(splitMaskEnv + strlen("OR"), &color)) ||
|
||||
(strncasecmp(splitMaskEnv, "|", strlen("|")) == 0 && parseInt(splitMaskEnv + strlen("|"), &color))
|
||||
)
|
||||
color = proc | color;
|
||||
if (
|
||||
(strncasecmp(splitMaskEnv, "MOD", strlen("MOD")) == 0 && parseInt(splitMaskEnv + strlen("MOD"), &color)) ||
|
||||
(strncasecmp(splitMaskEnv, "%", strlen("%")) == 0 && parseInt(splitMaskEnv + strlen("%"), &color))
|
||||
)
|
||||
color = proc % color;
|
||||
if (
|
||||
(strncasecmp(splitMaskEnv, "DIV", strlen("DIV")) == 0 && parseInt(splitMaskEnv + strlen("DIV"), &color)) ||
|
||||
(strncasecmp(splitMaskEnv, "/", strlen("/")) == 0 && parseInt(splitMaskEnv + strlen("/"), &color))
|
||||
)
|
||||
color = proc / color;
|
||||
}
|
||||
|
||||
MPI_Comm mpi_comm;
|
||||
color = proc & mask;
|
||||
MPI_Comm_split(MPI_COMM_WORLD, color, proc, &mpi_comm);
|
||||
MPI_Comm_size(mpi_comm, &ncclProcs);
|
||||
MPI_Comm_rank(mpi_comm, &ncclProc);
|
||||
@@ -1340,11 +1400,13 @@ testResult_t run() {
|
||||
int rank = proc*nThreads*nGpus+i;
|
||||
cudaDeviceProp prop;
|
||||
CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
|
||||
char busIdStr[] = "00000000:00:00.0";
|
||||
CUDACHECK(cudaDeviceGetPCIBusId(busIdStr, sizeof(busIdStr), cudaDev));
|
||||
len += snprintf(line+len, MAX_LINE>len ? MAX_LINE-len : 0, "# Rank %2d Pid %6d on %10s device %2d [%s] %s\n",
|
||||
rank, getpid(), hostname, cudaDev, busIdStr, prop.name);
|
||||
maxMem = std::min(maxMem, prop.totalGlobalMem);
|
||||
//char busIdStr[] = "00000000:00:00.0";
|
||||
//CUDACHECK(cudaDeviceGetPCIBusId(busIdStr, sizeof(busIdStr), cudaDev));
|
||||
//len += snprintf(line+len, MAX_LINE>len ? MAX_LINE-len : 0, "# Rank %2d Group %2d Pid %6d on %10s device %2d [%04x:%s:%02x] %s\n",
|
||||
// rank, color, getpid(), hostname, cudaDev, prop.pciDomainID, busIdStr, prop.pciDeviceID, prop.name);
|
||||
len += snprintf(line+len, MAX_LINE>len ? MAX_LINE-len : 0, "# Rank %2d Group %2d Pid %6d on %10s device %2d [%04x:%02x:%02x] %s\n",
|
||||
rank, color, getpid(), hostname, cudaDev, prop.pciDomainID, prop.pciBusID, prop.pciDeviceID, prop.name);
|
||||
maxMem = std::min(maxMem, prop.totalGlobalMem);
|
||||
}
|
||||
#if MPI_SUPPORT
|
||||
char *lines = (proc == 0) ? (char *)malloc(totalProcs*MAX_LINE) : NULL;
|
||||
@@ -1376,11 +1438,11 @@ testResult_t run() {
|
||||
MPI_Barrier(MPI_COMM_WORLD); // Ensure Bcast is complete for HCOLL
|
||||
#endif
|
||||
|
||||
int gpus[nGpus*nThreads];
|
||||
cudaStream_t streams[nGpus*nThreads];
|
||||
void* sendbuffs[nGpus*nThreads];
|
||||
void* recvbuffs[nGpus*nThreads];
|
||||
void* expected[nGpus*nThreads];
|
||||
std::vector<int> gpus(nGpus*nThreads);
|
||||
std::vector<cudaStream_t> streams(nGpus*nThreads);
|
||||
std::vector<void*> sendbuffs(nGpus*nThreads);
|
||||
std::vector<void*> recvbuffs(nGpus*nThreads);
|
||||
std::vector<void*> expected(nGpus*nThreads);
|
||||
size_t sendBytes, recvBytes;
|
||||
|
||||
ncclTestEngine.getBuffSize(&sendBytes, &recvBytes, (size_t)maxBytes, (size_t)ncclProcs*nGpus*nThreads);
|
||||
@@ -1390,11 +1452,11 @@ testResult_t run() {
|
||||
for (int i=0; i<nGpus*nThreads; i++) {
|
||||
gpus[i] = ((gpu0 != -1 ? gpu0 : localRank*nThreads*nGpus) + i)%numDevices;
|
||||
CUDACHECK(cudaSetDevice(gpus[i]));
|
||||
TESTCHECK(AllocateBuffs(sendbuffs+i, sendBytes, recvbuffs+i, recvBytes, expected+i, (size_t)maxBytes));
|
||||
TESTCHECK(AllocateBuffs(sendbuffs.data()+i, sendBytes, recvbuffs.data()+i, recvBytes, expected.data()+i, (size_t)maxBytes));
|
||||
if (streamnull)
|
||||
streams[i] = NULL;
|
||||
else
|
||||
CUDACHECK(cudaStreamCreateWithFlags(streams+i, cudaStreamNonBlocking));
|
||||
CUDACHECK(cudaStreamCreateWithFlags(streams.data()+i, cudaStreamNonBlocking));
|
||||
}
|
||||
|
||||
//if parallel init is not selected, use main thread to initialize NCCL
|
||||
@@ -1405,7 +1467,7 @@ testResult_t run() {
|
||||
#endif
|
||||
if (!parallel_init) {
|
||||
if (ncclProcs == 1) {
|
||||
NCCLCHECK(ncclCommInitAll(comms, nGpus*nThreads, gpus));
|
||||
NCCLCHECK(ncclCommInitAll(comms, nGpus*nThreads, gpus.data()));
|
||||
} else {
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
for (int i=0; i<nGpus*nThreads; i++) {
|
||||
@@ -1418,17 +1480,17 @@ testResult_t run() {
|
||||
sendRegHandles = (local_register) ? (void **)malloc(sizeof(*sendRegHandles)*nThreads*nGpus) : NULL;
|
||||
recvRegHandles = (local_register) ? (void **)malloc(sizeof(*recvRegHandles)*nThreads*nGpus) : NULL;
|
||||
for (int i=0; i<nGpus*nThreads; i++) {
|
||||
if (local_register) NCCLCHECK(ncclCommRegister(comms[i], sendbuffs[i], sendBytes, &sendRegHandles[i]));
|
||||
if (local_register) NCCLCHECK(ncclCommRegister(comms[i], recvbuffs[i], recvBytes, &recvRegHandles[i]));
|
||||
if (local_register) NCCLCHECK(ncclCommRegister(comms[i], &sendbuffs[i], maxBytes, &sendRegHandles[i]));
|
||||
if (local_register) NCCLCHECK(ncclCommRegister(comms[i], &recvbuffs[i], maxBytes, &recvRegHandles[i]));
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
int errors[nThreads];
|
||||
double bw[nThreads];
|
||||
std::vector<int> errors(nThreads);
|
||||
std::vector<double> bw(nThreads);
|
||||
double* delta;
|
||||
CUDACHECK(hipHostMalloc(&delta, sizeof(double)*nThreads*NUM_BLOCKS, cudaHostAllocPortable | cudaHostAllocMapped));
|
||||
int bw_count[nThreads];
|
||||
std::vector<int> bw_count(nThreads);
|
||||
for (int t=0; t<nThreads; t++) {
|
||||
bw[t] = 0.0;
|
||||
errors[t] = bw_count[t] = 0;
|
||||
@@ -1453,8 +1515,8 @@ testResult_t run() {
|
||||
}
|
||||
Reporter reporter(output_file, output_format);
|
||||
|
||||
struct testThread threads[nThreads];
|
||||
memset(threads, 0, sizeof(struct testThread)*nThreads);
|
||||
std::vector<testThread> threads(nThreads);
|
||||
memset(threads.data(), 0, sizeof(struct testThread)*nThreads);
|
||||
|
||||
for (int t=nThreads-1; t>=0; t--) {
|
||||
threads[t].args.minbytes=minBytes;
|
||||
@@ -1469,26 +1531,26 @@ testResult_t run() {
|
||||
threads[t].args.nThreads=nThreads;
|
||||
threads[t].args.thread=t;
|
||||
threads[t].args.nGpus=nGpus;
|
||||
threads[t].args.gpus=gpus+t*nGpus;
|
||||
threads[t].args.sendbuffs = sendbuffs+t*nGpus;
|
||||
threads[t].args.recvbuffs = recvbuffs+t*nGpus;
|
||||
threads[t].args.expected = expected+t*nGpus;
|
||||
threads[t].args.gpus=gpus.data()+t*nGpus;
|
||||
threads[t].args.sendbuffs = sendbuffs.data()+t*nGpus;
|
||||
threads[t].args.recvbuffs = recvbuffs.data()+t*nGpus;
|
||||
threads[t].args.expected = expected.data()+t*nGpus;
|
||||
threads[t].args.ncclId = ncclId;
|
||||
threads[t].args.comms=comms+t*nGpus;
|
||||
threads[t].args.streams=streams+t*nGpus;
|
||||
threads[t].args.streams=streams.data()+t*nGpus;
|
||||
threads[t].args.enable_out_of_place=enable_out_of_place;
|
||||
threads[t].args.enable_cache_flush = enable_cache_flush;
|
||||
threads[t].args.enable_rotating_tensor = enable_rotating_tensor;
|
||||
threads[t].args.errors=errors+t;
|
||||
threads[t].args.bw=bw+t;
|
||||
threads[t].args.bw_count=bw_count+t;
|
||||
threads[t].args.errors=errors.data()+t;
|
||||
threads[t].args.bw=bw.data()+t;
|
||||
threads[t].args.bw_count=bw_count.data()+t;
|
||||
|
||||
threads[t].args.reportErrors = datacheck;
|
||||
threads[t].args.reporter = &reporter;
|
||||
|
||||
threads[t].func = parallel_init ? threadInit : threadRunTests;
|
||||
if (t)
|
||||
TESTCHECK(threadLaunch(threads+t));
|
||||
TESTCHECK(threadLaunch(threads.data()+t));
|
||||
else
|
||||
TESTCHECK(threads[t].func(&threads[t].args));
|
||||
}
|
||||
|
||||
@@ -92,7 +92,7 @@ struct testColl {
|
||||
void (*getCollByteCount)(
|
||||
size_t *sendcount, size_t *recvcount, size_t *paramcount,
|
||||
size_t *sendInplaceOffset, size_t *recvInplaceOffset,
|
||||
size_t count, int nranks);
|
||||
size_t count, size_t eltSize, int nranks);
|
||||
testResult_t (*initData)(struct threadArgs* args, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, int rep, int in_place);
|
||||
void (*getBw)(size_t count, int typesize, double sec, double* algBw, double* busBw, int nranks);
|
||||
@@ -110,7 +110,7 @@ class Reporter {
|
||||
public:
|
||||
Reporter(std::string fileName, std::string outputFormat);
|
||||
~Reporter() { if (_outputValid) { _out.close(); } };
|
||||
void setParameters(const char* name, const char* typeName, const char* opName);// {
|
||||
void setParameters(const size_t numCycle, const char* name, const char* typeName, const char* opName);
|
||||
void addResult(int gpusPerRank, int ranksPerNode, int totalRanks, size_t numBytes, int inPlace, double timeUsec, double algBw, double busBw, int64_t wrongElts = -1);
|
||||
|
||||
private:
|
||||
@@ -121,6 +121,7 @@ class Reporter {
|
||||
bool _outputValid = false;
|
||||
std::ofstream _out;
|
||||
std::string _outputFormat;
|
||||
size_t _numCycle = 0;
|
||||
std::string _collectiveName;
|
||||
std::string _typeName;
|
||||
std::string _opName;
|
||||
@@ -294,7 +295,7 @@ extern ncclRedOp_t test_ops[];
|
||||
extern const char *test_opnames[];
|
||||
|
||||
static int ncclstringtotype(char *str) {
|
||||
for (int t=0; t<ncclNumTypes; t++) {
|
||||
for (int t=0; t<test_typenum; t++) {
|
||||
if (strcmp(str, test_typenames[t]) == 0) {
|
||||
return t;
|
||||
}
|
||||
|
||||
@@ -8,12 +8,12 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
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;
|
||||
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 = count/nranks;
|
||||
*paramcount = *sendcount;
|
||||
}
|
||||
|
||||
testResult_t GatherInitData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int rep, int in_place) {
|
||||
@@ -74,7 +74,7 @@ struct testColl gatherTest = {
|
||||
|
||||
void GatherGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
GatherGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
GatherGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t GatherRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -10,8 +10,8 @@
|
||||
|
||||
#define ALIGN 4
|
||||
|
||||
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;
|
||||
void HyperCubeGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) {
|
||||
size_t base = (count/nranks) & -(16/eltSize);
|
||||
*sendcount = base;
|
||||
*recvcount = base*nranks;
|
||||
*sendInplaceOffset = base;
|
||||
@@ -78,7 +78,7 @@ struct testColl hyperCubeTest = {
|
||||
|
||||
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);
|
||||
HyperCubeGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t HyperCubeRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -8,7 +8,7 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
void ReduceGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) {
|
||||
void ReduceGetCollByteCount(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;
|
||||
*recvcount = count;
|
||||
*sendInplaceOffset = 0;
|
||||
@@ -55,7 +55,7 @@ struct testColl reduceTest = {
|
||||
|
||||
void ReduceGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
ReduceGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
ReduceGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t ReduceRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -8,10 +8,8 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
#define ALIGN 4
|
||||
|
||||
void ReduceScatterGetCollByteCount(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;
|
||||
void ReduceScatterGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) {
|
||||
size_t base = (count/nranks) & -(16/eltSize);
|
||||
*sendcount = base*nranks;
|
||||
*recvcount = base;
|
||||
*sendInplaceOffset = 0;
|
||||
@@ -60,7 +58,7 @@ struct testColl reduceScatterTest = {
|
||||
|
||||
void ReduceScatterGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
ReduceScatterGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
ReduceScatterGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t ReduceScatterRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -8,12 +8,12 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
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;
|
||||
void ScatterGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) {
|
||||
*recvcount = (count/nranks) & -(16/eltSize);
|
||||
*sendcount = (*recvcount)*nranks;
|
||||
*sendInplaceOffset = 0;
|
||||
*recvInplaceOffset = count/nranks;
|
||||
*paramcount = count/nranks;
|
||||
*recvInplaceOffset = *recvcount;
|
||||
*paramcount = *recvcount;
|
||||
}
|
||||
|
||||
testResult_t ScatterInitData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int rep, int in_place) {
|
||||
@@ -70,7 +70,7 @@ struct testColl scatterTest = {
|
||||
|
||||
void ScatterGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
ScatterGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
ScatterGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t ScatterRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
@@ -8,7 +8,7 @@
|
||||
#include "cuda_runtime.h"
|
||||
#include "common.h"
|
||||
|
||||
void SendRecvGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, int nranks) {
|
||||
void SendRecvGetCollByteCount(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;
|
||||
*recvcount = count;
|
||||
*sendInplaceOffset = 0;
|
||||
@@ -69,7 +69,7 @@ struct testColl sendRecvTest = {
|
||||
|
||||
void SendRecvGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
SendRecvGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, nranks);
|
||||
SendRecvGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset, &recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t SendRecvRunTest(struct threadArgs* args, int root, ncclDataType_t type, const char* typeName, ncclRedOp_t op, const char* opName) {
|
||||
|
||||
새 이슈에서 참조
사용자 차단