Update RCCL-tests in examples folder (#261)
- Create a local copy for ROCm/rccl-tests for our examples.
- Update argument parsing to no longer use getopt_long.
- Workaround for Dyninst instrumentation.
---------
Co-authored-by: David Galiffi <David.Galiffi@amd.com>
[ROCm/rocprofiler-systems commit: 4e5029221b]
Cette révision appartient à :
Pranjal Swarup
GitHub
@@ -1,3 +1,25 @@
|
||||
# MIT License
|
||||
#
|
||||
# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved.
|
||||
#
|
||||
# Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
# of this software and associated documentation files (the "Software"), to deal
|
||||
# in the Software without restriction, including without limitation the rights
|
||||
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
# copies of the Software, and to permit persons to whom the Software is
|
||||
# furnished to do so, subject to the following conditions:
|
||||
#
|
||||
# The above copyright notice and this permission notice shall be included in all
|
||||
# copies or substantial portions of the Software.
|
||||
#
|
||||
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
|
||||
# SOFTWARE.
|
||||
|
||||
cmake_minimum_required(VERSION 3.18.4 FATAL_ERROR)
|
||||
|
||||
project(rocprofiler-systems-rccl-example LANGUAGES CXX)
|
||||
@@ -28,7 +50,7 @@ function(rccl_message _MSG_TYPE)
|
||||
endif()
|
||||
endfunction()
|
||||
|
||||
find_package(hip HINTS ${ROCmVersion_DIR} PATHS ${ROCmVersion_DIR})
|
||||
find_package(hip HINTS ${ROCmVersion_DIR} ${ROCM_PATH} /opt/rocm)
|
||||
|
||||
if(NOT hip_FOUND)
|
||||
rccl_message(AUTHOR_WARNING "${PROJECT_NAME} skipped. Missing HIP...")
|
||||
@@ -39,9 +61,10 @@ if(
|
||||
"${CMAKE_PROJECT_NAME}" STREQUAL "rocprofiler-systems"
|
||||
AND ("$ENV{ROCPROFSYS_CI}" OR ROCPROFSYS_CI OR ROCPROFSYS_BUILD_CI)
|
||||
)
|
||||
find_package(rccl QUIET) # avoid generating warning in CI
|
||||
# avoid generating warning in CI
|
||||
find_package(rccl HINTS ${ROCmVersion_DIR} ${ROCM_PATH} /opt/rocm QUIET)
|
||||
else()
|
||||
find_package(rccl)
|
||||
find_package(rccl HINTS ${ROCmVersion_DIR} ${ROCM_PATH} /opt/rocm)
|
||||
endif()
|
||||
|
||||
if(NOT rccl_FOUND)
|
||||
@@ -50,19 +73,22 @@ if(NOT rccl_FOUND)
|
||||
endif()
|
||||
|
||||
if(hip_FOUND AND rccl_FOUND)
|
||||
include(FetchContent)
|
||||
FetchContent_Declare(rccl-tests GIT_REPOSITORY https://github.com/ROCm/rccl-tests.git)
|
||||
# Use existing rccl-tests source in the project binary directory
|
||||
set(rccl-tests_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/rccl-tests")
|
||||
set(rccl-tests_BUILD_DIR "${CMAKE_CURRENT_BINARY_DIR}/rccl-tests")
|
||||
file(COPY ${rccl-tests_SOURCE_DIR}/ DESTINATION ${rccl-tests_BUILD_DIR})
|
||||
|
||||
# After the following call, the CMake targets defined by googletest and Catch2 will be
|
||||
# available to the rest of the build
|
||||
FetchContent_Populate(rccl-tests)
|
||||
if(NOT EXISTS "${rccl-tests_SOURCE_DIR}")
|
||||
rccl_message(AUTHOR_WARNING "${PROJECT_NAME} skipped. Missing rccl-tests source at ${rccl-tests_SOURCE_DIR}")
|
||||
return()
|
||||
endif()
|
||||
|
||||
get_filename_component(rccl_ROOT_DIR "${rccl_INCLUDE_DIR}" DIRECTORY)
|
||||
|
||||
rccl_message(STATUS "Building rccl-tests...")
|
||||
execute_process(
|
||||
COMMAND make HIP_HOME=${ROCM_PATH} RCCL_HOME=${rccl_ROOT_DIR}
|
||||
WORKING_DIRECTORY ${CMAKE_BINARY_DIR}/_deps/rccl-tests-src
|
||||
WORKING_DIRECTORY ${rccl-tests_BUILD_DIR}
|
||||
RESULT_VARIABLE _RCCL_BUILD_RET
|
||||
ERROR_VARIABLE _RCCL_BUILD_ERR
|
||||
OUTPUT_VARIABLE _RCCL_BUILD_OUT
|
||||
@@ -74,10 +100,7 @@ if(hip_FOUND AND rccl_FOUND)
|
||||
rccl_message(${_RCCL_BUILD_OUT})
|
||||
rccl_message(AUTHOR_WARNING "Failed to build rccl-tests: ${_RCCL_BUILD_ERR}")
|
||||
else()
|
||||
file(
|
||||
GLOB RCCL_TEST_EXECUTABLES
|
||||
${CMAKE_BINARY_DIR}/_deps/rccl-tests-src/build/*_perf
|
||||
)
|
||||
file(GLOB RCCL_TEST_EXECUTABLES ${rccl-tests_BUILD_DIR}/build/*_perf)
|
||||
set(_RCCL_TEST_TARGETS)
|
||||
|
||||
foreach(_EXE ${RCCL_TEST_EXECUTABLES})
|
||||
|
||||
@@ -0,0 +1,23 @@
|
||||
#
|
||||
# Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
|
||||
#
|
||||
# See LICENCE.txt for license information
|
||||
#
|
||||
|
||||
BUILDDIR ?= build
|
||||
override BUILDDIR := $(abspath $(BUILDDIR))
|
||||
|
||||
.PHONY: all clean
|
||||
|
||||
default: src.build
|
||||
|
||||
TARGETS=src
|
||||
|
||||
all: ${TARGETS:%=%.build}
|
||||
clean: ${TARGETS:%=%.clean}
|
||||
|
||||
%.build:
|
||||
${MAKE} -C $* build BUILDDIR=${BUILDDIR}
|
||||
|
||||
%.clean:
|
||||
${MAKE} -C $* clean BUILDDIR=${BUILDDIR}
|
||||
@@ -0,0 +1,174 @@
|
||||
#
|
||||
# Copyright (c) 2015-2022, NVIDIA CORPORATION. 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/x86_64-linux-gnu
|
||||
PREFIX ?= /usr/local
|
||||
VERBOSE ?= 0
|
||||
DEBUG ?= 0
|
||||
NCCL_HOME ?= ""
|
||||
CUSTOM_RCCL_LIB ?= ""
|
||||
|
||||
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
|
||||
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
|
||||
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
|
||||
else
|
||||
HIPCUFLAGS += -O0 -g -ggdb3
|
||||
endif
|
||||
|
||||
ifeq ($(VERBOSE), 0)
|
||||
.SILENT:
|
||||
endif
|
||||
|
||||
.PHONY: build clean
|
||||
|
||||
BUILDDIR ?= ../build
|
||||
HIPIFY_DIR ?= $(BUILDDIR)/hipify
|
||||
|
||||
.PRECIOUS: $(HIPIFY_DIR)/%.cpp $(HIPIFY_DIR)/%.h
|
||||
|
||||
ifeq ($(MPI), 1)
|
||||
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${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%)
|
||||
|
||||
DST_DIR := $(BUILDDIR)
|
||||
SRC_FILES := $(wildcard *.cpp)
|
||||
OBJ_FILES := $(SRC_FILES:%.cpp=${DST_DIR}/%.o)
|
||||
BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall scatter gather sendrecv alltoallv
|
||||
BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf)
|
||||
|
||||
GIT_VERSION_FILE := ${DST_DIR}/src/git_version.cpp
|
||||
GIT_REV := $(shell git log --pretty=format:'%h' -n 1)
|
||||
GIT_DIFF := $(shell git diff --quiet --exit-code || echo +)
|
||||
GIT_BRANCH := $(shell git rev-parse --abbrev-ref HEAD)
|
||||
|
||||
build: ${BIN_FILES}
|
||||
|
||||
clean:
|
||||
rm -rf ${DST_DIR}
|
||||
|
||||
TEST_VERIFIABLE_SRCDIR := ../verifiable
|
||||
TEST_VERIFIABLE_BUILDDIR := $(BUILDDIR)/verifiable
|
||||
include ../verifiable/verifiable.mk
|
||||
|
||||
# Rule to create git_version.cpp
|
||||
$(GIT_VERSION_FILE):
|
||||
@mkdir -p ${DST_DIR}/src
|
||||
@echo 'const char* rcclTestsGitHash = "$(GIT_BRANCH):$(GIT_REV)$(GIT_DIFF)";' > $@
|
||||
|
||||
${HIPIFY_DIR}/%.cpp: %.cpp
|
||||
@printf "Hipifying %-35s > %s\n" $< $@
|
||||
@mkdir -p ${HIPIFY_DIR}
|
||||
hipify-perl -quiet-warnings $< > $@
|
||||
|
||||
${HIPIFY_DIR}/%.h: %.h
|
||||
@printf "Hipifying %-35s > %s\n" $< $@
|
||||
@mkdir -p ${HIPIFY_DIR}
|
||||
hipify-perl -quiet-warnings $< > $@
|
||||
|
||||
${DST_DIR}/%.o: ${HIPIFY_DIR}/%.cpp ${HIPIFY_DIR}/common.h $(TEST_VERIFIABLE_HDRS) $(GIT_VERSION_FILE)
|
||||
@printf "Compiling %-35s > %s\n" $< $@
|
||||
@mkdir -p ${DST_DIR}
|
||||
echo "$(HIPCC) $(HIPCUFLAGS) -I. -c -o $@ $<"
|
||||
$(HIPCC) $(HIPCUFLAGS) -I. -c -o $@ $<
|
||||
|
||||
${DST_DIR}/timer.o: timer.cc timer.h
|
||||
@printf "Compiling %-35s > %s\n" $< $@
|
||||
@mkdir -p ${DST_DIR}
|
||||
$(CXX) $(CXXFLAGS) -o $@ -c timer.cc
|
||||
|
||||
${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 $@ $^ $(HIPLDFLAGS)"
|
||||
$(HIPCC) -o $@ $^ $(HIPLDFLAGS)
|
||||
|
||||
@@ -0,0 +1,110 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = base;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AllGatherInitData(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++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? ((char*) args->recvbuffs[i]) + rank * args->sendBytes
|
||||
: args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, ncclSum, 33 * rep + rank, 1, 0));
|
||||
for(int j = 0; j < nranks; j++)
|
||||
{
|
||||
TESTCHECK(InitData((char*) args->expected[i] + args->sendBytes * j, sendcount,
|
||||
0, type, ncclSum, 33 * rep + j, 1, 0));
|
||||
}
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
AllGatherGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * typesize * nranks) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double) (nranks - 1)) / ((double) nranks);
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AllGatherRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
NCCLCHECK(ncclAllGather(sendbuff, recvbuff, count, type, comm, stream));
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl allGatherTest = { "AllGather", AllGatherGetCollByteCount,
|
||||
AllGatherInitData, AllGatherGetBw, AllGatherRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AllGatherRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &allGatherTest;
|
||||
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, "none", -1));
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { AllGatherGetBuffSize, AllGatherRunTest };
|
||||
@@ -0,0 +1,126 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = *sendcount;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AllReduceInitData(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++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, op, rep, nranks, rank));
|
||||
TESTCHECK(InitDataReduce(args->expected[i], recvcount, 0, type, op, rep, nranks));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
AllReduceGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double) (2 * (nranks - 1))) / ((double) nranks);
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AllReduceRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
NCCLCHECK(ncclAllReduce(sendbuff, recvbuff, count, type, op, comm, stream));
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl allReduceTest = { "AllReduce", AllReduceGetCollByteCount,
|
||||
AllReduceInitData, AllReduceGetBw, AllReduceRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AllReduceRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &allReduceTest;
|
||||
ncclDataType_t* run_types;
|
||||
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 = 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++)
|
||||
{
|
||||
for(int j = 0; j < op_count; j++)
|
||||
{
|
||||
#if defined(RCCL_FLOAT8)
|
||||
if((run_types[i] == ncclFp8E4M3 || run_types[i] == ncclFp8E5M2) &&
|
||||
run_ops[j] == ncclProd)
|
||||
continue;
|
||||
#endif
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], run_ops[j],
|
||||
run_opnames[j], -1));
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { AllReduceGetBuffSize, AllReduceRunTest };
|
||||
@@ -0,0 +1,112 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AlltoAllInitData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root,
|
||||
int rep, int in_place)
|
||||
{
|
||||
size_t sendcount = args->sendBytes / wordSize(type);
|
||||
size_t recvcount = args->expectedBytes / wordSize(type);
|
||||
int nranks = args->nProcs * args->nThreads * args->nGpus;
|
||||
|
||||
for(int i = 0; i < args->nGpus; i++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, ncclSum, 33 * rep + rank, 1, 0));
|
||||
for(int j = 0; j < nranks; j++)
|
||||
{
|
||||
size_t partcount = sendcount / nranks;
|
||||
TESTCHECK(InitData((char*) args->expected[i] + j * partcount * wordSize(type),
|
||||
partcount, rank * partcount, type, ncclSum, 33 * rep + j,
|
||||
1, 0));
|
||||
}
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
// We don't support in-place alltoall
|
||||
args->reportErrors = in_place ? 0 : 1;
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
AlltoAllGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * nranks * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double) (nranks - 1)) / ((double) (nranks));
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AlltoAllRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
NCCLCHECK(ncclAllToAll(sendbuff, recvbuff, count, type, comm, stream));
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl alltoAllTest = { "AlltoAll", AlltoAllGetCollByteCount, AlltoAllInitData,
|
||||
AlltoAllGetBw, AlltoAllRunColl };
|
||||
|
||||
void
|
||||
AlltoAllGetBuffSize(size_t* sendcount, size_t* recvcount, size_t count, int nranks)
|
||||
{
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
AlltoAllGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset,
|
||||
&recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AlltoAllRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &alltoAllTest;
|
||||
ncclDataType_t* run_types;
|
||||
const char** run_typenames;
|
||||
int type_count;
|
||||
|
||||
if((int) type != -1)
|
||||
{
|
||||
type_count = 1;
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
}
|
||||
else
|
||||
{
|
||||
type_count = 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, "none", -1));
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { AlltoAllGetBuffSize, AlltoAllRunTest };
|
||||
@@ -0,0 +1,218 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2020, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
#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, size_t eltSize, int nranks)
|
||||
{
|
||||
if(count < nranks * nranks / 2)
|
||||
{
|
||||
*sendcount = 0;
|
||||
*recvcount = 0;
|
||||
*sendInplaceOffset = 0;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = 0;
|
||||
}
|
||||
else
|
||||
{
|
||||
*paramcount = (count / nranks) & -(16 / eltSize);
|
||||
*sendcount = nranks * (*paramcount);
|
||||
*recvcount = *sendcount;
|
||||
*sendInplaceOffset = 0;
|
||||
*recvInplaceOffset = 0;
|
||||
}
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AlltoAllvInitData(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++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, ncclSum, 33 * rep + rank, 1, 0));
|
||||
|
||||
#if 0
|
||||
int *dataHost = (int *)malloc(args->sendBytes);
|
||||
cudaMemcpy(dataHost, data, args->sendBytes, cudaMemcpyDeviceToHost);
|
||||
printf(" Rank [%d] Original: ", rank);
|
||||
for(int j=0; j<sendcount; j++) {
|
||||
printf("%d:%d ", j, dataHost[j]);
|
||||
}
|
||||
printf("\n");
|
||||
free(dataHost);
|
||||
#endif
|
||||
|
||||
size_t rdisp = 0;
|
||||
size_t data_count = sendcount * 2 / nranks;
|
||||
size_t chunksize = data_count / nranks;
|
||||
for(int j = 0; j < nranks; j++)
|
||||
{
|
||||
size_t scount = 0, rcount = ((j + rank) % nranks) * chunksize;
|
||||
if((j + rank) % nranks == 0)
|
||||
rcount += (sendcount - chunksize * (nranks - 1) * nranks / 2);
|
||||
size_t sdisp = 0;
|
||||
for(int k = 0; k < nranks; k++)
|
||||
{
|
||||
scount = ((k + j) % nranks) * chunksize;
|
||||
if((k + j) % nranks == 0)
|
||||
scount += (sendcount - chunksize * (nranks - 1) * nranks / 2);
|
||||
if(k == rank) break;
|
||||
sdisp += scount;
|
||||
}
|
||||
TESTCHECK(InitData(((char*) args->expected[i]) + rdisp * wordSize(type),
|
||||
rcount, sdisp, type, ncclSum, 33 * rep + j, 1, 0));
|
||||
rdisp += rcount;
|
||||
}
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
// We don't support in-place alltoall
|
||||
args->reportErrors = in_place ? 0 : 1;
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
AlltoAllvGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * nranks * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double) (nranks - 1)) / ((double) (nranks));
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AlltoAllvRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
int nranks;
|
||||
NCCLCHECK(ncclCommCount(comm, &nranks));
|
||||
int rank;
|
||||
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
||||
|
||||
if(count == 0) return testSuccess;
|
||||
|
||||
size_t *sendcounts, *recvcounts, *sdispls, *rdispls;
|
||||
sendcounts = (size_t*) malloc(nranks * nranks * sizeof(size_t));
|
||||
recvcounts = (size_t*) malloc(nranks * nranks * sizeof(size_t));
|
||||
sdispls = (size_t*) malloc(nranks * nranks * sizeof(size_t));
|
||||
rdispls = (size_t*) malloc(nranks * nranks * sizeof(size_t));
|
||||
if(sendcounts == nullptr || recvcounts == nullptr || sdispls == nullptr ||
|
||||
rdispls == nullptr)
|
||||
{
|
||||
printf("failed to allocate buffers for alltoallv\n");
|
||||
return testNcclError;
|
||||
}
|
||||
|
||||
size_t disp = 0;
|
||||
size_t chunksize = count * 2 / nranks;
|
||||
for(int i = 0; i < nranks; i++)
|
||||
{
|
||||
size_t scount = ((i + rank) % nranks) * chunksize;
|
||||
if((i + rank) % nranks == 0)
|
||||
scount += (count * nranks - chunksize * (nranks - 1) * nranks / 2);
|
||||
sendcounts[i + rank * nranks] = recvcounts[i + rank * nranks] = scount;
|
||||
sdispls[i + rank * nranks] = rdispls[i + rank * nranks] = disp;
|
||||
disp += scount;
|
||||
// printf("%d->%d: sendcounts/recvcounts %lx sdispls/rdispls %lx\n", rank, i,
|
||||
// sendcounts[i+rank*nranks]*wordSize(type),
|
||||
// sdispls[i+rank*nranks]*wordSize(type));
|
||||
}
|
||||
|
||||
#if NCCL_MAJOR < 2 || NCCL_MINOR < 7
|
||||
printf(
|
||||
"NCCL 2.7 or later is needed for alltoallv. This test was compiled with %d.%d.\n",
|
||||
NCCL_MAJOR, NCCL_MINOR);
|
||||
return testNcclError;
|
||||
#else
|
||||
# if defined(RCCL_ALLTOALLV) && defined(USE_RCCL_GATHER_SCATTER)
|
||||
NCCLCHECK(ncclAllToAllv(sendbuff, sendcounts + rank * nranks, sdispls + rank * nranks,
|
||||
recvbuff, recvcounts + rank * nranks, rdispls + rank * nranks,
|
||||
type, comm, stream));
|
||||
# else
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
for(int r = 0; r < nranks; r++)
|
||||
{
|
||||
if(sendcounts[r + rank * nranks] != 0)
|
||||
{
|
||||
NCCLCHECK(
|
||||
ncclSend(((char*) sendbuff) + sdispls[r + rank * nranks] * wordSize(type),
|
||||
sendcounts[r + rank * nranks], type, r, comm, stream));
|
||||
}
|
||||
if(recvcounts[r + rank * nranks] != 0)
|
||||
{
|
||||
NCCLCHECK(
|
||||
ncclRecv(((char*) recvbuff) + rdispls[r + rank * nranks] * wordSize(type),
|
||||
recvcounts[r + rank * nranks], type, r, comm, stream));
|
||||
}
|
||||
}
|
||||
NCCLCHECK(ncclGroupEnd());
|
||||
# endif
|
||||
#endif
|
||||
free(sendcounts);
|
||||
free(recvcounts);
|
||||
free(sdispls);
|
||||
free(rdispls);
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl alltoAllTest = { "AlltoAllv", AlltoAllvGetCollByteCount,
|
||||
AlltoAllvInitData, AlltoAllvGetBw, AlltoAllvRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
AlltoAllvRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &alltoAllTest;
|
||||
ncclDataType_t* run_types;
|
||||
const char** run_typenames;
|
||||
int type_count;
|
||||
|
||||
if((int) type != -1)
|
||||
{
|
||||
type_count = 1;
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
}
|
||||
else
|
||||
{
|
||||
type_count = ncclNumTypes;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
|
||||
for(int i = 0; i < type_count; i++)
|
||||
{
|
||||
TESTCHECK(
|
||||
TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t) 0, "", -1));
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { AlltoAllvGetBuffSize, AlltoAllvRunTest };
|
||||
@@ -0,0 +1,131 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = *sendcount;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
BroadcastInitData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root,
|
||||
int rep, int in_place)
|
||||
{
|
||||
size_t sendcount = args->sendBytes / wordSize(type);
|
||||
size_t recvcount = args->expectedBytes / wordSize(type);
|
||||
|
||||
for(int i = 0; i < args->nGpus; i++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
if(rank == root)
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, ncclSum, rep, 1, 0));
|
||||
TESTCHECK(InitData(args->expected[i], recvcount, 0, type, ncclSum, rep, 1, 0));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
BroadcastGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = 1;
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
BroadcastRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
int rank;
|
||||
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
||||
#if NCCL_MAJOR >= 2 && NCCL_MINOR >= 2
|
||||
NCCLCHECK(ncclBroadcast(sendbuff, recvbuff, count, type, root, comm, stream));
|
||||
#else
|
||||
if(rank == root)
|
||||
{
|
||||
NCCLCHECK(ncclBcast(sendbuff, count, type, root, comm, stream));
|
||||
}
|
||||
else
|
||||
{
|
||||
NCCLCHECK(ncclBcast(recvbuff, count, type, root, comm, stream));
|
||||
}
|
||||
#endif
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl broadcastTest = { "Broadcast", BroadcastGetCollByteCount,
|
||||
BroadcastInitData, BroadcastGetBw, BroadcastRunColl };
|
||||
|
||||
void
|
||||
BroadcastGetBuffSize(size_t* sendcount, size_t* recvcount, size_t count, int nranks)
|
||||
{
|
||||
size_t paramcount, sendInplaceOffset, recvInplaceOffset;
|
||||
BroadcastGetCollByteCount(sendcount, recvcount, ¶mcount, &sendInplaceOffset,
|
||||
&recvInplaceOffset, count, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
BroadcastRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &broadcastTest;
|
||||
ncclDataType_t* run_types;
|
||||
const char** run_typenames;
|
||||
int type_count;
|
||||
int begin_root, end_root;
|
||||
|
||||
if((int) type != -1)
|
||||
{
|
||||
type_count = 1;
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
}
|
||||
else
|
||||
{
|
||||
type_count = test_typenum;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
|
||||
if(root != -1)
|
||||
{
|
||||
begin_root = end_root = root;
|
||||
}
|
||||
else
|
||||
{
|
||||
begin_root = 0;
|
||||
end_root = args->nProcs * args->nThreads * args->nGpus - 1;
|
||||
}
|
||||
|
||||
for(int i = 0; i < type_count; i++)
|
||||
{
|
||||
for(int j = begin_root; j <= end_root; j++)
|
||||
{
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t) 0,
|
||||
"none", j));
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { BroadcastGetBuffSize, BroadcastRunTest };
|
||||
Fichier diff supprimé car celui-ci est trop grand
Voir la Diff
@@ -0,0 +1,431 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
#ifndef __COMMON_H__
|
||||
#define __COMMON_H__
|
||||
|
||||
#include "rccl/rccl.h"
|
||||
#include <algorithm>
|
||||
#include <cstdint>
|
||||
#include <stdio.h>
|
||||
#ifdef MPI_SUPPORT
|
||||
# include "mpi.h"
|
||||
#endif
|
||||
#include "nccl1_compat.h"
|
||||
#include "timer.h"
|
||||
#include <fstream>
|
||||
#include <iostream>
|
||||
#include <pthread.h>
|
||||
#include <string>
|
||||
|
||||
// Ensures backward compatibility for FP8 types in RCCL 2.24.3 and later
|
||||
#if NCCL_VERSION_CODE >= NCCL_VERSION(2, 24, 3)
|
||||
# define ncclFp8E4M3 ncclFloat8e4m3
|
||||
# define ncclFp8E5M2 ncclFloat8e5m2
|
||||
#endif
|
||||
|
||||
// For nccl.h < 2.13 since we define a weak fallback
|
||||
extern "C" char const*
|
||||
ncclGetLastError(ncclComm_t comm);
|
||||
|
||||
#define CUDACHECK(cmd) \
|
||||
do \
|
||||
{ \
|
||||
cudaError_t err = cmd; \
|
||||
if(err != cudaSuccess) \
|
||||
{ \
|
||||
char hostname[1024]; \
|
||||
getHostName(hostname, 1024); \
|
||||
printf("%s: Test CUDA failure %s:%d '%s'\n", hostname, __FILE__, __LINE__, \
|
||||
cudaGetErrorString(err)); \
|
||||
return testCudaError; \
|
||||
} \
|
||||
} while(0)
|
||||
|
||||
#if NCCL_VERSION_CODE >= NCCL_VERSION(2, 13, 0)
|
||||
# define NCCLCHECK(cmd) \
|
||||
do \
|
||||
{ \
|
||||
ncclResult_t res = cmd; \
|
||||
if(res != ncclSuccess) \
|
||||
{ \
|
||||
char hostname[1024]; \
|
||||
getHostName(hostname, 1024); \
|
||||
printf("%s: Test NCCL failure %s:%d " \
|
||||
"'%s / %s'\n", \
|
||||
hostname, __FILE__, __LINE__, ncclGetErrorString(res), \
|
||||
ncclGetLastError(NULL)); \
|
||||
return testNcclError; \
|
||||
} \
|
||||
} while(0)
|
||||
#else
|
||||
# define NCCLCHECK(cmd) \
|
||||
do \
|
||||
{ \
|
||||
ncclResult_t res = cmd; \
|
||||
if(res != ncclSuccess) \
|
||||
{ \
|
||||
char hostname[1024]; \
|
||||
getHostName(hostname, 1024); \
|
||||
printf("%s: Test NCCL failure %s:%d '%s'\n", hostname, __FILE__, \
|
||||
__LINE__, ncclGetErrorString(res)); \
|
||||
return testNcclError; \
|
||||
} \
|
||||
} while(0)
|
||||
#endif
|
||||
|
||||
typedef enum
|
||||
{
|
||||
testSuccess = 0,
|
||||
testInternalError = 1,
|
||||
testCudaError = 2,
|
||||
testNcclError = 3,
|
||||
testTimeout = 4,
|
||||
testNumResults = 5
|
||||
} testResult_t;
|
||||
|
||||
// Relay errors up and trace
|
||||
#define TESTCHECK(cmd) \
|
||||
do \
|
||||
{ \
|
||||
testResult_t r = cmd; \
|
||||
if(r != testSuccess) \
|
||||
{ \
|
||||
char hostname[1024]; \
|
||||
getHostName(hostname, 1024); \
|
||||
printf(" .. %s pid %d: Test failure %s:%d\n", hostname, getpid(), __FILE__, \
|
||||
__LINE__); \
|
||||
return r; \
|
||||
} \
|
||||
} while(0)
|
||||
|
||||
struct testColl
|
||||
{
|
||||
const char name[20];
|
||||
void (*getCollByteCount)(size_t* sendcount, size_t* recvcount, size_t* paramcount,
|
||||
size_t* sendInplaceOffset, size_t* recvInplaceOffset,
|
||||
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);
|
||||
testResult_t (*runColl)(void* sendbuff, void* recvbuff, size_t count,
|
||||
ncclDataType_t type, ncclRedOp_t op, int root,
|
||||
ncclComm_t comm, cudaStream_t stream);
|
||||
};
|
||||
extern struct testColl allReduceTest;
|
||||
extern struct testColl allGatherTest;
|
||||
extern struct testColl reduceScatterTest;
|
||||
extern struct testColl broadcastTest;
|
||||
extern struct testColl reduceTest;
|
||||
extern struct testColl alltoAllTest;
|
||||
|
||||
class Reporter
|
||||
{
|
||||
public:
|
||||
Reporter(std::string fileName, std::string outputFormat);
|
||||
~Reporter()
|
||||
{
|
||||
if(_outputValid)
|
||||
{
|
||||
_out.close();
|
||||
}
|
||||
};
|
||||
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:
|
||||
bool isMainThread();
|
||||
template <typename T>
|
||||
std::pair<std::string, std::string> makeValueKeyPair(T v, std::string k)
|
||||
{
|
||||
return std::make_pair(std::to_string(v), k);
|
||||
};
|
||||
template <>
|
||||
std::pair<std::string, std::string> makeValueKeyPair<std::string>(std::string v,
|
||||
std::string k)
|
||||
{
|
||||
return std::make_pair("\"" + v + "\"", k);
|
||||
};
|
||||
|
||||
bool _outputValid = false;
|
||||
std::ofstream _out;
|
||||
std::string _outputFormat;
|
||||
size_t _numCycle = 0;
|
||||
std::string _collectiveName;
|
||||
std::string _typeName;
|
||||
std::string _opName;
|
||||
};
|
||||
|
||||
struct testEngine
|
||||
{
|
||||
void (*getBuffSize)(size_t* sendcount, size_t* recvcount, size_t count, int nranks);
|
||||
testResult_t (*runTest)(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName);
|
||||
};
|
||||
|
||||
extern struct testEngine ncclTestEngine;
|
||||
|
||||
struct threadArgs
|
||||
{
|
||||
size_t nbytes;
|
||||
size_t minbytes;
|
||||
size_t maxbytes;
|
||||
size_t stepbytes;
|
||||
size_t stepfactor;
|
||||
|
||||
int totalProcs;
|
||||
int nProcs;
|
||||
int proc;
|
||||
int nThreads;
|
||||
int thread;
|
||||
int nGpus;
|
||||
int* gpus;
|
||||
int localRank;
|
||||
int enable_out_of_place;
|
||||
int enable_in_place;
|
||||
int enable_cache_flush;
|
||||
int enable_rotating_tensor;
|
||||
void** sendbuffs;
|
||||
size_t sendBytes;
|
||||
size_t sendInplaceOffset;
|
||||
void** recvbuffs;
|
||||
size_t recvInplaceOffset;
|
||||
ncclUniqueId ncclId;
|
||||
ncclComm_t* comms;
|
||||
cudaStream_t* streams;
|
||||
|
||||
void** expected;
|
||||
size_t expectedBytes;
|
||||
int* errors;
|
||||
double* bw;
|
||||
int* bw_count;
|
||||
|
||||
int reportErrors;
|
||||
|
||||
struct testColl* collTest;
|
||||
|
||||
Reporter* reporter;
|
||||
};
|
||||
|
||||
typedef testResult_t (*threadFunc_t)(struct threadArgs* args);
|
||||
struct testThread
|
||||
{
|
||||
pthread_t thread;
|
||||
threadFunc_t func;
|
||||
struct threadArgs args;
|
||||
testResult_t ret;
|
||||
};
|
||||
|
||||
// Provided by common.cu
|
||||
extern void
|
||||
Barrier(struct threadArgs* args);
|
||||
extern testResult_t
|
||||
TimeTest(struct threadArgs* args, ncclDataType_t type, const char* typeName,
|
||||
ncclRedOp_t op, const char* opName, int root);
|
||||
extern testResult_t
|
||||
InitDataReduce(void* data, const size_t count, const size_t offset, ncclDataType_t type,
|
||||
ncclRedOp_t op, const uint64_t seed, const int nranks);
|
||||
extern testResult_t
|
||||
InitData(void* data, const size_t count, size_t offset, ncclDataType_t type,
|
||||
ncclRedOp_t op, const uint64_t seed, const int nranks, const int rank);
|
||||
extern void
|
||||
AllocateBuffs(void** sendbuff, void** recvbuff, void** expected, void** expectedHost,
|
||||
size_t nbytes, int nranks);
|
||||
|
||||
#include <unistd.h>
|
||||
|
||||
static void
|
||||
getHostName(char* hostname, int maxlen)
|
||||
{
|
||||
gethostname(hostname, maxlen);
|
||||
for(int i = 0; i < maxlen; i++)
|
||||
{
|
||||
if(hostname[i] == '.')
|
||||
{
|
||||
hostname[i] = '\0';
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#include <stdint.h>
|
||||
|
||||
static uint64_t
|
||||
getHash(const char* string, size_t n)
|
||||
{
|
||||
// Based on DJB2a, result = result * 33 ^ char
|
||||
uint64_t result = 5381;
|
||||
for(size_t c = 0; c < n; c++)
|
||||
{
|
||||
result = ((result << 5) + result) ^ string[c];
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
/* Generate a hash of the unique identifying string for this host
|
||||
* that will be unique for both bare-metal and container instances
|
||||
* Equivalent of a hash of;
|
||||
*
|
||||
* $(hostname)$(cat /proc/sys/kernel/random/boot_id)
|
||||
*
|
||||
*/
|
||||
#define HOSTID_FILE "/proc/sys/kernel/random/boot_id"
|
||||
static uint64_t
|
||||
getHostHash(const char* hostname)
|
||||
{
|
||||
char hostHash[1024];
|
||||
|
||||
// Fall back is the hostname if something fails
|
||||
(void) strncpy(hostHash, hostname, sizeof(hostHash));
|
||||
int offset = strlen(hostHash);
|
||||
|
||||
FILE* file = fopen(HOSTID_FILE, "r");
|
||||
if(file != NULL)
|
||||
{
|
||||
char* p;
|
||||
if(fscanf(file, "%ms", &p) == 1)
|
||||
{
|
||||
strncpy(hostHash + offset, p, sizeof(hostHash) - offset - 1);
|
||||
free(p);
|
||||
}
|
||||
}
|
||||
fclose(file);
|
||||
|
||||
// Make sure the string is terminated
|
||||
hostHash[sizeof(hostHash) - 1] = '\0';
|
||||
|
||||
return getHash(hostHash, strlen(hostHash));
|
||||
}
|
||||
|
||||
static size_t
|
||||
wordSize(ncclDataType_t type)
|
||||
{
|
||||
switch(type)
|
||||
{
|
||||
case ncclChar:
|
||||
#if NCCL_MAJOR >= 2
|
||||
// case ncclInt8:
|
||||
case ncclUint8:
|
||||
# if NCCL_MAJOR >= 2 && RCCL_FLOAT8 == 1
|
||||
case ncclFp8E4M3:
|
||||
case ncclFp8E5M2:
|
||||
# endif
|
||||
#endif
|
||||
return 1;
|
||||
case ncclHalf:
|
||||
#if NCCL_MAJOR >= 2 && RCCL_BFLOAT16 == 1
|
||||
case ncclBfloat16:
|
||||
#endif
|
||||
// case ncclFloat16:
|
||||
return 2;
|
||||
case ncclInt:
|
||||
case ncclFloat:
|
||||
#if NCCL_MAJOR >= 2
|
||||
// case ncclInt32:
|
||||
case ncclUint32:
|
||||
// case ncclFloat32:
|
||||
#endif
|
||||
return 4;
|
||||
case ncclInt64:
|
||||
case ncclUint64:
|
||||
case ncclDouble:
|
||||
// case ncclFloat64:
|
||||
return 8;
|
||||
default: return 0;
|
||||
}
|
||||
}
|
||||
|
||||
extern int test_ncclVersion; // init'd with ncclGetVersion()
|
||||
typedef enum
|
||||
{
|
||||
ncclCoarse = 0,
|
||||
ncclFine = 1,
|
||||
ncclHost = 2,
|
||||
ncclManaged = 3,
|
||||
nccl_NUM_MTYPES = 4
|
||||
} ncclMemoryType_t;
|
||||
extern const char* test_memorytypes[nccl_NUM_MTYPES];
|
||||
constexpr int test_opNumMax =
|
||||
(int) ncclNumOps + (NCCL_VERSION_CODE >= NCCL_VERSION(2, 11, 0) ? 1 : 0);
|
||||
extern int test_opnum;
|
||||
extern int test_typenum;
|
||||
extern ncclDataType_t test_types[ncclNumTypes];
|
||||
extern const char* test_typenames[ncclNumTypes];
|
||||
extern ncclRedOp_t test_ops[];
|
||||
extern const char* test_opnames[];
|
||||
|
||||
static int
|
||||
ncclstringtotype(char* str)
|
||||
{
|
||||
for(int t = 0; t < test_typenum; t++)
|
||||
{
|
||||
if(strcmp(str, test_typenames[t]) == 0)
|
||||
{
|
||||
return t;
|
||||
}
|
||||
}
|
||||
if(strcmp(str, "all") == 0)
|
||||
{
|
||||
return -1;
|
||||
}
|
||||
printf("invalid type %s, defaulting to %s .. \n", str, test_typenames[ncclFloat]);
|
||||
return ncclFloat;
|
||||
}
|
||||
|
||||
static int
|
||||
ncclstringtoop(char* str)
|
||||
{
|
||||
for(int o = 0; o < test_opnum; o++)
|
||||
{
|
||||
if(strcmp(str, test_opnames[o]) == 0)
|
||||
{
|
||||
return o;
|
||||
}
|
||||
}
|
||||
if(strcmp(str, "all") == 0)
|
||||
{
|
||||
return -1;
|
||||
}
|
||||
printf("invalid op %s, defaulting to %s .. \n", str, test_opnames[ncclSum]);
|
||||
return ncclSum;
|
||||
}
|
||||
|
||||
static int
|
||||
ncclstringtoroot(char* str)
|
||||
{
|
||||
if(strcmp(str, "all") == 0)
|
||||
{
|
||||
return -1;
|
||||
}
|
||||
return strtol(str, NULL, 0);
|
||||
}
|
||||
|
||||
static int
|
||||
ncclstringtomtype(char* str)
|
||||
{
|
||||
for(int o = 0; o < nccl_NUM_MTYPES; o++)
|
||||
{
|
||||
if(strcmp(str, test_memorytypes[o]) == 0)
|
||||
{
|
||||
return o;
|
||||
}
|
||||
}
|
||||
printf("invalid memorytype %s, defaulting to %s .. \n", str,
|
||||
test_memorytypes[ncclCoarse]);
|
||||
return ncclCoarse;
|
||||
}
|
||||
|
||||
extern int is_main_proc;
|
||||
extern thread_local int is_main_thread;
|
||||
#define PRINT \
|
||||
if(is_main_thread) printf
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,143 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
void
|
||||
GatherGetCollByteCount(size_t* sendcount, size_t* recvcount, size_t* paramcount,
|
||||
size_t* sendInplaceOffset, size_t* recvInplaceOffset, size_t count,
|
||||
size_t eltSize, int nranks)
|
||||
{
|
||||
*sendcount = (count / nranks) & -(16 / eltSize);
|
||||
*recvcount = (*sendcount) * nranks;
|
||||
*sendInplaceOffset = *sendcount;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = *sendcount;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
GatherInitData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root,
|
||||
int rep, int in_place)
|
||||
{
|
||||
size_t sendcount = args->sendBytes / wordSize(type);
|
||||
size_t recvcount = args->expectedBytes / wordSize(type);
|
||||
int nranks = args->nProcs * args->nThreads * args->nGpus;
|
||||
|
||||
for(int i = 0; i < args->nGpus; i++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? ((char*) args->recvbuffs[i]) + rank * args->sendBytes
|
||||
: args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, rank * sendcount, type, ncclSum, rep, 1, 0));
|
||||
CUDACHECK(cudaMemcpy(args->expected[i], args->recvbuffs[i], args->expectedBytes,
|
||||
cudaMemcpyDefault));
|
||||
if(rank == root)
|
||||
{
|
||||
TESTCHECK(InitData(args->expected[i], nranks * sendcount, 0, type, ncclSum,
|
||||
rep, 1, 0));
|
||||
}
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
GatherGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * nranks * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double) (nranks - 1)) / ((double) (nranks));
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
GatherRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
int nRanks;
|
||||
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
||||
int rank;
|
||||
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
||||
size_t rankOffset = count * wordSize(type);
|
||||
if(count == 0) return testSuccess;
|
||||
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
NCCLCHECK(ncclSend(sendbuff, count, type, root, comm, stream));
|
||||
if(rank == root)
|
||||
{
|
||||
for(int r = 0; r < nRanks; r++)
|
||||
{
|
||||
NCCLCHECK(ncclRecv(((char*) recvbuff) + r * rankOffset, count, type, r, comm,
|
||||
stream));
|
||||
}
|
||||
}
|
||||
NCCLCHECK(ncclGroupEnd());
|
||||
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl gatherTest = { "Gather", GatherGetCollByteCount, GatherInitData,
|
||||
GatherGetBw, GatherRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
GatherRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &gatherTest;
|
||||
ncclDataType_t* run_types;
|
||||
const char** run_typenames;
|
||||
int type_count;
|
||||
int begin_root, end_root;
|
||||
|
||||
if((int) type != -1)
|
||||
{
|
||||
type_count = 1;
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
}
|
||||
else
|
||||
{
|
||||
type_count = test_typenum;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
|
||||
if(root != -1)
|
||||
{
|
||||
begin_root = end_root = root;
|
||||
}
|
||||
else
|
||||
{
|
||||
begin_root = 0;
|
||||
end_root = args->nProcs * args->nThreads * args->nGpus - 1;
|
||||
}
|
||||
|
||||
for(int i = 0; i < type_count; i++)
|
||||
{
|
||||
for(int j = begin_root; j <= end_root; j++)
|
||||
{
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t) 0,
|
||||
"none", j));
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { GatherGetBuffSize, GatherRunTest };
|
||||
@@ -0,0 +1,6 @@
|
||||
#ifndef RCCL_TESTS_GIT_VERSION_H_
|
||||
#define RCCL_TESTS_GIT_VERSION_H_
|
||||
|
||||
extern const char* rcclTestsGitHash;
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,144 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
#define ALIGN 4
|
||||
|
||||
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;
|
||||
*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++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? ((char*) args->recvbuffs[i]) + rank * args->sendBytes
|
||||
: args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, ncclSum, 33 * rep + rank, 1, 0));
|
||||
for(int j = 0; j < nranks; j++)
|
||||
{
|
||||
TESTCHECK(InitData((char*) args->expected[i] + args->sendBytes * j, sendcount,
|
||||
0, type, ncclSum, 33 * rep + j, 1, 0));
|
||||
}
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
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, cudaStream_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)
|
||||
CUDACHECK(cudaMemcpyAsync(rbuff + rank * rankSize, sbuff, rankSize,
|
||||
cudaMemcpyDeviceToDevice, 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, /*eltSize=*/1, 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;
|
||||
}
|
||||
|
||||
// Check if this is a power of 2
|
||||
int nRanks = args->nProcs * args->nThreads * args->nGpus;
|
||||
if(nRanks && !(nRanks & (nRanks - 1)))
|
||||
{
|
||||
for(int i = 0; i < type_count; i++)
|
||||
{
|
||||
TESTCHECK(
|
||||
TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t) 0, "", -1));
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
printf("nRanks %d is not a power of 2, skipping\n", nRanks);
|
||||
}
|
||||
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { HyperCubeGetBuffSize, HyperCubeRunTest };
|
||||
@@ -0,0 +1,70 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#ifndef NCCL1_COMPAT_H
|
||||
#define NCCL1_COMPAT_H
|
||||
|
||||
#ifndef NCCL_MAJOR // NCCL 1.x
|
||||
# define NCCL_MAJOR 1
|
||||
# define NCCL_MINOR 0
|
||||
|
||||
# define ncclNumOps nccl_NUM_OPS
|
||||
# define ncclNumTypes nccl_NUM_TYPES
|
||||
|
||||
static ncclResult_t
|
||||
ncclGroupStart()
|
||||
{
|
||||
return ncclSuccess;
|
||||
}
|
||||
static ncclResult_t
|
||||
ncclGroupEnd()
|
||||
{
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
# define CHECKCOUNT(count) \
|
||||
if(count > INT_MAX) return ncclInvalidArgument;
|
||||
|
||||
static ncclResult_t
|
||||
ncclReduce(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
CHECKCOUNT(count);
|
||||
return ncclReduce(sendbuff, recvbuff, (int) count, datatype, op, root, comm, stream);
|
||||
}
|
||||
static ncclResult_t
|
||||
ncclAllReduce(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype,
|
||||
ncclRedOp_t op, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
CHECKCOUNT(count);
|
||||
return ncclAllReduce(sendbuff, recvbuff, (int) count, datatype, op, comm, stream);
|
||||
}
|
||||
static ncclResult_t
|
||||
ncclBcast(void* buff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm,
|
||||
cudaStream_t stream)
|
||||
{
|
||||
CHECKCOUNT(count);
|
||||
return ncclBcast(buff, (int) count, datatype, root, comm, stream);
|
||||
}
|
||||
static ncclResult_t
|
||||
ncclReduceScatter(const void* sendbuff, void* recvbuff, size_t recvcount,
|
||||
ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm,
|
||||
cudaStream_t stream)
|
||||
{
|
||||
CHECKCOUNT(recvcount);
|
||||
return ncclReduceScatter(sendbuff, recvbuff, (int) recvcount, datatype, op, comm,
|
||||
stream);
|
||||
}
|
||||
static ncclResult_t
|
||||
ncclAllGather(const void* sendbuff, void* recvbuff, size_t sendcount,
|
||||
ncclDataType_t datatype, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
CHECKCOUNT(sendcount);
|
||||
return ncclAllGather(sendbuff, (int) sendcount, datatype, recvbuff, comm, stream);
|
||||
}
|
||||
#endif
|
||||
|
||||
#endif
|
||||
Fichier diff supprimé car celui-ci est trop grand
Voir la Diff
@@ -0,0 +1,142 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = *sendcount;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ReduceInitData(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++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, op, rep, nranks, rank));
|
||||
CUDACHECK(cudaMemcpy(args->expected[i], args->recvbuffs[i], args->expectedBytes,
|
||||
cudaMemcpyDefault));
|
||||
if(rank == root)
|
||||
TESTCHECK(
|
||||
InitDataReduce(args->expected[i], recvcount, 0, type, op, rep, nranks));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
ReduceGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * typesize) / 1.0E9 / sec;
|
||||
*algBw = baseBw;
|
||||
*busBw = baseBw;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ReduceRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
NCCLCHECK(ncclReduce(sendbuff, recvbuff, count, type, op, root, comm, stream));
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl reduceTest = { "Reduce", ReduceGetCollByteCount, ReduceInitData,
|
||||
ReduceGetBw, ReduceRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ReduceRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &reduceTest;
|
||||
ncclDataType_t* run_types;
|
||||
ncclRedOp_t* run_ops;
|
||||
const char ** run_typenames, **run_opnames;
|
||||
int type_count, op_count;
|
||||
int begin_root, end_root;
|
||||
|
||||
if((int) type != -1)
|
||||
{
|
||||
type_count = 1;
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
}
|
||||
else
|
||||
{
|
||||
type_count = test_typenum;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
|
||||
if((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;
|
||||
}
|
||||
|
||||
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++)
|
||||
{
|
||||
for(int j = 0; j < op_count; j++)
|
||||
{
|
||||
#if defined(RCCL_FLOAT8)
|
||||
if((run_types[i] == ncclFp8E4M3 || run_types[i] == ncclFp8E5M2) &&
|
||||
run_ops[j] == ncclProd)
|
||||
continue;
|
||||
#endif
|
||||
for(int k = begin_root; k <= end_root; k++)
|
||||
{
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], run_ops[j],
|
||||
run_opnames[j], k));
|
||||
}
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { ReduceGetBuffSize, ReduceRunTest };
|
||||
@@ -0,0 +1,131 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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;
|
||||
*recvInplaceOffset = base;
|
||||
*paramcount = base;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ReduceScatterInitData(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++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, op, rep, nranks, rank));
|
||||
CUDACHECK(cudaMemcpy(args->expected[i], args->recvbuffs[i], args->expectedBytes,
|
||||
cudaMemcpyDefault));
|
||||
TESTCHECK(InitDataReduce(args->expected[i], recvcount, rank * recvcount, type, op,
|
||||
rep, nranks));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
ReduceScatterGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * typesize * nranks) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double) (nranks - 1)) / ((double) nranks);
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ReduceScatterRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
NCCLCHECK(ncclReduceScatter(sendbuff, recvbuff, count, type, op, comm, stream));
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl reduceScatterTest = { "ReduceScatter", ReduceScatterGetCollByteCount,
|
||||
ReduceScatterInitData, ReduceScatterGetBw,
|
||||
ReduceScatterRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ReduceScatterRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &reduceScatterTest;
|
||||
ncclDataType_t* run_types;
|
||||
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 = test_typenum;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
|
||||
if((int) op != -1)
|
||||
{
|
||||
run_ops = &op;
|
||||
run_opnames = &opName;
|
||||
op_count = 1;
|
||||
}
|
||||
else
|
||||
{
|
||||
op_count = test_opnum;
|
||||
run_ops = test_ops;
|
||||
run_opnames = test_opnames;
|
||||
}
|
||||
|
||||
for(int i = 0; i < type_count; i++)
|
||||
{
|
||||
for(int j = 0; j < op_count; j++)
|
||||
{
|
||||
#if defined(RCCL_FLOAT8)
|
||||
if((run_types[i] == ncclFp8E4M3 || run_types[i] == ncclFp8E5M2) &&
|
||||
run_ops[j] == ncclProd)
|
||||
continue;
|
||||
#endif
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], run_ops[j],
|
||||
run_opnames[j], -1));
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { ReduceScatterGetBuffSize, ReduceScatterRunTest };
|
||||
@@ -0,0 +1,137 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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 = *recvcount;
|
||||
*paramcount = *recvcount;
|
||||
}
|
||||
|
||||
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);
|
||||
|
||||
for(int i = 0; i < args->nGpus; i++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
if(rank == root)
|
||||
TESTCHECK(InitData(data, sendcount, 0, type, ncclSum, rep, 1, 0));
|
||||
TESTCHECK(InitData(args->expected[i], recvcount, rank * recvcount, type, ncclSum,
|
||||
rep, 1, 0));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
ScatterGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * nranks * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = ((double) (nranks - 1)) / ((double) (nranks));
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ScatterRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_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;
|
||||
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
if(rank == root)
|
||||
{
|
||||
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());
|
||||
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testColl scatterTest = { "Scatter", ScatterGetCollByteCount, ScatterInitData,
|
||||
ScatterGetBw, ScatterRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
ScatterRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &scatterTest;
|
||||
ncclDataType_t* run_types;
|
||||
const char** run_typenames;
|
||||
int type_count;
|
||||
int begin_root, end_root;
|
||||
|
||||
if((int) type != -1)
|
||||
{
|
||||
type_count = 1;
|
||||
run_types = &type;
|
||||
run_typenames = &typeName;
|
||||
}
|
||||
else
|
||||
{
|
||||
type_count = test_typenum;
|
||||
run_types = test_types;
|
||||
run_typenames = test_typenames;
|
||||
}
|
||||
|
||||
if(root != -1)
|
||||
{
|
||||
begin_root = end_root = root;
|
||||
}
|
||||
else
|
||||
{
|
||||
begin_root = 0;
|
||||
end_root = args->nProcs * args->nThreads * args->nGpus - 1;
|
||||
}
|
||||
|
||||
for(int i = 0; i < type_count; i++)
|
||||
{
|
||||
for(int j = begin_root; j <= end_root; j++)
|
||||
{
|
||||
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], (ncclRedOp_t) 0,
|
||||
"none", j));
|
||||
}
|
||||
}
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { ScatterGetBuffSize, ScatterRunTest };
|
||||
@@ -0,0 +1,135 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#include "common.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
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;
|
||||
*recvInplaceOffset = 0;
|
||||
*paramcount = *sendcount;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
SendRecvInitData(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++)
|
||||
{
|
||||
CUDACHECK(cudaSetDevice(args->gpus[i]));
|
||||
int rank = ((args->proc * args->nThreads + args->thread) * args->nGpus + i);
|
||||
CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
|
||||
void* data = in_place ? args->recvbuffs[i] : args->sendbuffs[i];
|
||||
TESTCHECK(InitData(data, sendcount, rank * sendcount, type, ncclSum, rep, 1, 0));
|
||||
int peer = (rank - 1 + nranks) % nranks;
|
||||
TESTCHECK(InitData(args->expected[i], recvcount, peer * recvcount, type, ncclSum,
|
||||
rep, 1, 0));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
// We don't support in-place sendrecv
|
||||
args->reportErrors = in_place ? 0 : 1;
|
||||
return testSuccess;
|
||||
}
|
||||
|
||||
void
|
||||
SendRecvGetBw(size_t count, int typesize, double sec, double* algBw, double* busBw,
|
||||
int nranks)
|
||||
{
|
||||
double baseBw = (double) (count * typesize) / 1.0E9 / sec;
|
||||
|
||||
*algBw = baseBw;
|
||||
double factor = 1;
|
||||
*busBw = baseBw * factor;
|
||||
}
|
||||
|
||||
testResult_t
|
||||
SendRecvRunColl(void* sendbuff, void* recvbuff, size_t count, ncclDataType_t type,
|
||||
ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream)
|
||||
{
|
||||
int nRanks;
|
||||
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
||||
int rank;
|
||||
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
||||
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 = { "SendRecv", SendRecvGetCollByteCount, SendRecvInitData,
|
||||
SendRecvGetBw, SendRecvRunColl };
|
||||
|
||||
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, /*eltSize=*/1, nranks);
|
||||
}
|
||||
|
||||
testResult_t
|
||||
SendRecvRunTest(struct threadArgs* args, int root, ncclDataType_t type,
|
||||
const char* typeName, ncclRedOp_t op, const char* opName)
|
||||
{
|
||||
args->collTest = &sendRecvTest;
|
||||
ncclDataType_t* run_types;
|
||||
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 = 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++)
|
||||
{
|
||||
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;
|
||||
}
|
||||
|
||||
struct testEngine ncclTestEngine = { SendRecvGetBuffSize, SendRecvRunTest };
|
||||
@@ -0,0 +1,35 @@
|
||||
#include "timer.h"
|
||||
|
||||
// Make sure to compile this translation unit with the host compiler and not
|
||||
// nvcc, lest you hit an internal compiler error (ICE) with GCC 10.3.0
|
||||
#include <chrono>
|
||||
|
||||
namespace
|
||||
{
|
||||
std::uint64_t
|
||||
now()
|
||||
{
|
||||
using clock = std::chrono::steady_clock;
|
||||
return std::chrono::duration_cast<std::chrono::nanoseconds>(
|
||||
clock::now().time_since_epoch())
|
||||
.count();
|
||||
}
|
||||
} // namespace
|
||||
|
||||
timer::timer() { t0 = now(); }
|
||||
|
||||
double
|
||||
timer::elapsed() const
|
||||
{
|
||||
std::uint64_t t1 = now();
|
||||
return 1.e-9 * (t1 - t0);
|
||||
}
|
||||
|
||||
double
|
||||
timer::reset()
|
||||
{
|
||||
std::uint64_t t1 = now();
|
||||
double ans = 1.e-9 * (t1 - t0);
|
||||
t0 = t1;
|
||||
return ans;
|
||||
}
|
||||
@@ -0,0 +1,17 @@
|
||||
#ifndef _408319ecdd5b47b28bf8f511c4fdf816
|
||||
#define _408319ecdd5b47b28bf8f511c4fdf816
|
||||
|
||||
#include <cstdint>
|
||||
|
||||
// Can't include <chrono> because of bug with gcc 10.3.0
|
||||
class timer
|
||||
{
|
||||
std::uint64_t t0;
|
||||
|
||||
public:
|
||||
timer();
|
||||
double elapsed() const;
|
||||
double reset();
|
||||
};
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,74 @@
|
||||
#
|
||||
# Copyright (c) 2015-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
# Modifications are Copyright (c) 2019-2025 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# See LICENSE.txt for license information
|
||||
#
|
||||
|
||||
#include ../../makefiles/common.mk
|
||||
|
||||
.PHONY: all clean
|
||||
|
||||
BUILDDIR := $(abspath ../../build)
|
||||
DST_DIR := $(BUILDDIR)/test/verifiable
|
||||
|
||||
ROCM_PATH ?= /opt/rocm
|
||||
MPI_HOME ?= /usr/lib/x86_64-linux-gnu
|
||||
PREFIX ?= /usr/local
|
||||
VERBOSE ?= 0
|
||||
DEBUG ?= 0
|
||||
NCCL_HOME ?= ""
|
||||
|
||||
HIPCC = $(ROCM_PATH)/bin/amdclang++
|
||||
CXX = $(HIPCC)
|
||||
|
||||
HIPCUFLAGS := -std=c++14
|
||||
LDFLAGS :=
|
||||
HIPLDFLAGS :=
|
||||
|
||||
ifneq ($(NCCL_HOME), "")
|
||||
HIPCUFLAGS += -I$(NCCL_HOME)/ -I$(NCCL_HOME)/include
|
||||
HIPLDFLAGS += -Wl,-rpath,$(NCCL_HOME) -L$(NCCL_HOME) -L$(NCCL_HOME)/lib
|
||||
endif
|
||||
|
||||
HIPCUFLAGS += -I$(ROCM_PATH)/include
|
||||
HIPCUFLAGS += -I$(ROCM_PATH)/include/hip
|
||||
LDFLAGS += -L$(ROCM_PATH)/lib -lhsa-runtime64 -lrt
|
||||
HIPLDFLAGS += $(CUSTOM_RCCL_LIB) -L$(ROCM_PATH)/lib -lhsa-runtime64 -lamdhip64 -lstdc++ -lrt
|
||||
|
||||
ifeq ($(DEBUG), 0)
|
||||
HIPCUFLAGS += -O3
|
||||
else
|
||||
HIPCUFLAGS += -O0 -g -ggdb3
|
||||
endif
|
||||
|
||||
ifeq ($(VERBOSE), 0)
|
||||
.SILENT:
|
||||
endif
|
||||
|
||||
ifeq ($(MPI), 1)
|
||||
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${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%)
|
||||
|
||||
all: $(DST_DIR)/verifiable.o $(DST_DIR)/self_test
|
||||
|
||||
clean:
|
||||
rm -rf $(DST_DIR)
|
||||
|
||||
TEST_VERIFIABLE_SRCDIR := .
|
||||
TEST_VERIFIABLE_BUILDDIR := $(DST_DIR)
|
||||
include verifiable.mk
|
||||
|
||||
self_test: $(DST_DIR)/self_test
|
||||
|
||||
$(DST_DIR)/self_test: verifiable.cpp verifiable.h
|
||||
@printf "Linking %s\n" $@
|
||||
@mkdir -p $(DST_DIR)
|
||||
$(HIPCC) -o $@ $(HIPCUFLAGS) -DSELF_TEST=1 verifiable.cpp $(HIPLDFLAGS)
|
||||
+236
@@ -0,0 +1,236 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
/* Generate parameters for our error bound model of floating point average
|
||||
* (sum of scaled values) by sampling sums of random sequences for each
|
||||
* floating point type.
|
||||
*
|
||||
* The model has parameters "coef" and "power", where for two floats a & b,
|
||||
* they are close enough if and only if:
|
||||
* abs(intBits(a) - intBits(b)) <= 1 + coef*pow(rank_n, power);
|
||||
*
|
||||
* Where intBits(x) is the reinterpretation of the float bitpattern as an integer.
|
||||
*
|
||||
* Compile with:
|
||||
* nvcc -gencode=arch=compute_80,code=sm_80
|
||||
*/
|
||||
|
||||
#include <algorithm>
|
||||
#include <cmath>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cuda_fp16.h>
|
||||
#include <hip/hip_bfloat16.h>
|
||||
|
||||
using std::uint32_t;
|
||||
using std::uint64_t;
|
||||
using bfloat16 = hip_bfloat16;
|
||||
|
||||
template <typename T>
|
||||
struct float_traits;
|
||||
|
||||
template <>
|
||||
struct float_traits<float>
|
||||
{
|
||||
static constexpr int mantissa_bits = 23;
|
||||
static constexpr int exponent_bits = 8;
|
||||
using uint_t = uint32_t;
|
||||
__device__ static float make(double x) { return (float) x; }
|
||||
__device__ static float make(uint64_t x) { return (float) x; }
|
||||
__device__ static double todouble(float x) { return x; }
|
||||
__device__ static float add(float a, float b) { return a + b; }
|
||||
__device__ static float mul(float a, float b) { return a * b; }
|
||||
};
|
||||
template <>
|
||||
struct float_traits<double>
|
||||
{
|
||||
static constexpr int mantissa_bits = 52;
|
||||
static constexpr int exponent_bits = 11;
|
||||
using uint_t = uint64_t;
|
||||
__device__ static double make(double x) { return x; }
|
||||
__device__ static double make(uint64_t x) { return (double) x; }
|
||||
__device__ static double todouble(double x) { return x; }
|
||||
__device__ static double add(double a, double b) { return a + b; }
|
||||
__device__ static double mul(double a, double b) { return a * b; }
|
||||
};
|
||||
template <>
|
||||
struct float_traits<__half>
|
||||
{
|
||||
static constexpr int mantissa_bits = 10;
|
||||
static constexpr int exponent_bits = 5;
|
||||
using uint_t = uint16_t;
|
||||
__device__ static __half make(double x) { return __float2half((float) x); }
|
||||
__device__ static __half make(uint64_t x) { return __int2half_rn(x); }
|
||||
__device__ static double todouble(__half x) { return __half2float(x); }
|
||||
__device__ static __half add(__half a, __half b) { return __hadd(a, b); }
|
||||
__device__ static __half mul(__half a, __half b) { return __hmul(a, b); }
|
||||
};
|
||||
template <>
|
||||
struct float_traits<bfloat16>
|
||||
{
|
||||
static constexpr int mantissa_bits = 7;
|
||||
static constexpr int exponent_bits = 8;
|
||||
using uint_t = uint16_t;
|
||||
__device__ static bfloat16 make(double x) { return bfloat16(x); }
|
||||
__device__ static bfloat16 make(uint64_t x) { return bfloat16(x); }
|
||||
__device__ static double todouble(bfloat16 x) { return double(x); }
|
||||
__device__ static bfloat16 add(bfloat16 a, bfloat16 b)
|
||||
{
|
||||
return bfloat16(__hadd((float) a, (float) b));
|
||||
}
|
||||
__device__ static bfloat16 mul(bfloat16 a, bfloat16 b)
|
||||
{
|
||||
return bfloat16(__hmul((float) a, (float) b));
|
||||
}
|
||||
};
|
||||
|
||||
template <typename F>
|
||||
__device__ int
|
||||
compare(F a, F b)
|
||||
{
|
||||
union
|
||||
{
|
||||
typename float_traits<F>::uint_t ua;
|
||||
F fa;
|
||||
};
|
||||
union
|
||||
{
|
||||
typename float_traits<F>::uint_t ub;
|
||||
F fb;
|
||||
};
|
||||
ua = 0;
|
||||
ub = 0;
|
||||
fa = a;
|
||||
fb = b;
|
||||
// std::printf("bits(%1.10f)=%x bits(%1.10f)=%x\n", fa, ua, fb, ub);
|
||||
return ua < ub ? ub - ua : ua - ub;
|
||||
}
|
||||
|
||||
struct xoshiro256ss
|
||||
{
|
||||
uint64_t s[4];
|
||||
__device__ xoshiro256ss(int seed)
|
||||
{
|
||||
constexpr uint64_t src[4] = { 0xbb99e851d1f545cc, 0xbfc4022389ca40cb,
|
||||
0xe84aff5cb1914af5, 0x845999858284de77 };
|
||||
for(int i = 0; i < 4; i++)
|
||||
s[i] = src[i] + (seed + i) * 0xb45de8a52fdb65d3;
|
||||
}
|
||||
__device__ uint64_t operator()()
|
||||
{
|
||||
auto rol64 = [](uint64_t x, int k) { return (x << k) | (x >> (64 - k)); };
|
||||
uint64_t const result = rol64(s[1] * 5, 7) * 9;
|
||||
uint64_t const t = s[1] << 17;
|
||||
s[2] ^= s[0];
|
||||
s[3] ^= s[1];
|
||||
s[1] ^= s[2];
|
||||
s[0] ^= s[3];
|
||||
s[2] ^= t;
|
||||
s[3] = rol64(s[3], 45);
|
||||
return result;
|
||||
}
|
||||
};
|
||||
|
||||
static __device__ int
|
||||
__reduce_max_sync(unsigned int mask, int value)
|
||||
{
|
||||
// We ignore mask, since all bits are set when calling them in the
|
||||
// test code below.
|
||||
int width = warpSize;
|
||||
for(unsigned int i = warpSize; i; i >>= 1)
|
||||
{
|
||||
value = max(__shfl_down(value, i, width), value);
|
||||
}
|
||||
return value;
|
||||
}
|
||||
|
||||
template <typename F>
|
||||
__global__ void
|
||||
kernel()
|
||||
{
|
||||
using traits = float_traits<F>;
|
||||
constexpr int samps = 4 << 10;
|
||||
__shared__ F accf[samps];
|
||||
__shared__ double accd[samps];
|
||||
|
||||
xoshiro256ss rng(threadIdx.x);
|
||||
float expo_avg = 1;
|
||||
for(int pass = 0; pass < 2; pass++)
|
||||
{
|
||||
F scalar = traits::make(1.0 / (3.14159 + .5 * threadIdx.x));
|
||||
int err_max = 0;
|
||||
float coef = 0;
|
||||
double expo_sum = 0;
|
||||
int expo_n = 0;
|
||||
int max_ranks =
|
||||
std::is_same<F, float>::value ? 16 << 10 : 1 << traits::mantissa_bits;
|
||||
for(int round = 0; round < 1 + (16 << 10) / max_ranks; round++)
|
||||
{
|
||||
// for(int round=0; round < 2; round++) {
|
||||
for(int i = threadIdx.x; i < samps; i += blockDim.x)
|
||||
{
|
||||
accf[i] = (F) 0;
|
||||
accd[i] = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
for(int r = 0; r < max_ranks; r++)
|
||||
{
|
||||
int err = 0;
|
||||
for(int i = threadIdx.x; i < samps; i += blockDim.x)
|
||||
{
|
||||
constexpr uint64_t m = (1ll << traits::mantissa_bits) - 1;
|
||||
double d = std::is_same<F, float>::value ? double(rng() & m) : 1.0;
|
||||
F f = traits::make(d);
|
||||
accf[i] = traits::add(accf[i], traits::mul(scalar, f));
|
||||
accd[i] += traits::todouble(f);
|
||||
// if(threadIdx.x==0 && std::is_same<F,half>::value) std::printf("
|
||||
// r=%d f=%f\n", r, traits::todouble(accf[i]));
|
||||
int e = compare(accf[i], traits::mul(scalar, traits::make(accd[i])));
|
||||
err = err > e ? err : e;
|
||||
}
|
||||
err = __reduce_max_sync(-1u, err);
|
||||
err_max = err_max > err ? err_max : err;
|
||||
if(r >= 2)
|
||||
{
|
||||
// err = 1 + coef*pow(r,expo)
|
||||
float c = float(err - 1) / powf(float(r), expo_avg);
|
||||
coef = coef > c ? coef : c;
|
||||
}
|
||||
if(r >= 2)
|
||||
{
|
||||
double expo = log2f(1 + err_max) / log2f(r);
|
||||
expo_sum += expo;
|
||||
expo_n++;
|
||||
// if(threadIdx.x==0 && std::is_same<F,half>::value) std::printf("
|
||||
// r=%d err=%d errmax=%d expo=%f sum=%f n=%d\n", r, err, err_max,
|
||||
// expo, expo_sum, expo_n);
|
||||
}
|
||||
}
|
||||
}
|
||||
if(pass == 0)
|
||||
expo_avg = expo_sum / expo_n;
|
||||
else if(threadIdx.x == 0)
|
||||
printf(" coef=%1.10f expo=%1.10f\n", coef, expo_avg);
|
||||
}
|
||||
}
|
||||
|
||||
int
|
||||
main()
|
||||
{
|
||||
std::printf("type=float:\n");
|
||||
kernel<float><<<1, 32>>>();
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
std::printf("\ntype=half:\n");
|
||||
kernel<half><<<1, 32>>>();
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
std::printf("\ntype=bfloat16:\n");
|
||||
kernel<bfloat16><<<1, 32>>>();
|
||||
cudaDeviceSynchronize();
|
||||
return 0;
|
||||
}
|
||||
+1563
Fichier diff supprimé car celui-ci est trop grand
Voir la Diff
@@ -0,0 +1,67 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
* Modifications Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
|
||||
#ifndef _d41d8cd98f00b204e9800998ecf8427e
|
||||
#define _d41d8cd98f00b204e9800998ecf8427e
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <stdint.h>
|
||||
|
||||
/* Routines for launching kernels that verify reduction results. A significant
|
||||
* feature of these routines is they carefully craft floating point input
|
||||
* to produce exactly predictable output.
|
||||
*
|
||||
* int elt_ty: actually just a ncclDataType_t
|
||||
*
|
||||
* int red_op: mostly just a ncclRedOp_t. Since PreMulSum ops are dynamically
|
||||
* created, these are encoded as the value ncclNumOps and their scalar is
|
||||
* assumed to be `ncclVerifiablePremulScalar(rank_me)`
|
||||
*
|
||||
* uint64_t seed: arbitrary 64-bits to use in seeding the random values
|
||||
*
|
||||
* intptr_t elt_ix0: index of first element pointed to by elts when generating
|
||||
* random values. This makes it possible to generate subsequences independently
|
||||
* as well as in aggregate.
|
||||
*
|
||||
* int rank_n: Number of contributions into the reduction. Non-reduction
|
||||
* collectives like broadcast, gather, etc will always set this to one.
|
||||
*
|
||||
* int rank_me: Index of this contribution
|
||||
*/
|
||||
|
||||
// Use this as the local scalar for PreMulSum ops
|
||||
template <typename T>
|
||||
__host__ __device__ T
|
||||
ncclVerifiablePremulScalar(int rank_me)
|
||||
{
|
||||
return T(rank_me % 2 == 0 ? 1.0f : 2.0f);
|
||||
}
|
||||
|
||||
// Enqueue kernel to generate data which is to be reduced.
|
||||
void
|
||||
ncclVerifiablePrepareInput(void* elts, intptr_t elt_n, int elt_ty, int red_op, int rank_n,
|
||||
int rank_me, uint64_t seed, intptr_t elt_ix0,
|
||||
cudaStream_t stream);
|
||||
|
||||
// Enqueue kernel to generate expected results of reduction.
|
||||
void
|
||||
ncclVerifiablePrepareExpected(void* elts, intptr_t elt_n, int elt_ty, int red_op,
|
||||
int rank_n, uint64_t seed, intptr_t elt_ix0,
|
||||
cudaStream_t stream);
|
||||
|
||||
// Enqueue kernel to verify reduced data matches expectation. The number of
|
||||
// failed elements is written to bad_elt_n which must be in cudaHost memory.
|
||||
// If `expected == nullptr` then the expected results are generated on-the-fly
|
||||
// which can be costly. Thus if you plan to run the same reduction multiple
|
||||
// times it is advantageous to precompute the expected values with
|
||||
// ncclVerifiablePrepareExpected and pass them as `expected` here.
|
||||
void
|
||||
ncclVerifiableVerify(void const* results, void const* expected, intptr_t elt_n,
|
||||
int elt_ty, int red_op, int rank_n, uint64_t seed, intptr_t elt_ix0,
|
||||
int64_t* bad_elt_n, cudaStream_t stream);
|
||||
#endif
|
||||
@@ -0,0 +1,33 @@
|
||||
# Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
|
||||
# Modifications Copyright (c) 2020-2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# See LICENSE.txt for license information
|
||||
|
||||
|
||||
# We requires both of the following paths to be set upon including this makefile
|
||||
# TEST_VERIFIABLE_SRCDIR = <points to this directory>
|
||||
# TEST_VERIFIABLE_BUILDDIR = <points to destination of .o file>
|
||||
|
||||
TEST_VERIFIABLE_HDRS = $(TEST_VERIFIABLE_SRCDIR)/verifiable.h
|
||||
TEST_VERIFIABLE_OBJS = $(TEST_VERIFIABLE_BUILDDIR)/verifiable.o
|
||||
|
||||
${HIPIFY_DIR}/verifiable.cpp: $(TEST_VERIFIABLE_SRCDIR)/verifiable.cpp
|
||||
@printf "Hipifying %-35s > %s\n" $< $@
|
||||
@mkdir -p ${HIPIFY_DIR}
|
||||
hipify-perl -quiet-warnings $< > $@
|
||||
|
||||
${HIPIFY_DIR}/verifiable.h: $(TEST_VERIFIABLE_SRCDIR)/verifiable.h
|
||||
@printf "Hipifying %-35s > %s\n" $< $@
|
||||
@mkdir -p ${HIPIFY_DIR}
|
||||
hipify-perl -quiet-warnings $< > $@
|
||||
|
||||
${HIPIFY_DIR}/rccl_float8.h: $(TEST_VERIFIABLE_SRCDIR)/../src/rccl_float8.h
|
||||
@printf "Hipifying %-35s > %s\n" $< $@
|
||||
@mkdir -p ${HIPIFY_DIR}
|
||||
hipify-perl -quiet-warnings $< > $@
|
||||
|
||||
$(TEST_VERIFIABLE_BUILDDIR)/verifiable.o: $(HIPIFY_DIR)/verifiable.cpp $(HIPIFY_DIR)/verifiable.h $(HIPIFY_DIR)/rccl_float8.h
|
||||
@printf "Compiling %s\n" $@
|
||||
@mkdir -p $(TEST_VERIFIABLE_BUILDDIR)
|
||||
echo " $(HIPCC) -o $@ $(HIPCUFLAGS) -c $<"
|
||||
$(HIPCC) -o $@ $(HIPCUFLAGS) -c $<
|
||||
Référencer dans un nouveau ticket
Bloquer un utilisateur