Moved tests to separate dir and improved MPI test
test sources moved to test/ directory.
MPI test displays PASS/FAIL and returns code accordingly.
Change-Id: I058ebd1bd5202d8f38cc9787898b2480100c102b
Reviewed-on: http://git-master/r/936086
Reviewed-by: Przemek Tredak <ptredak@nvidia.com>
Tested-by: Przemek Tredak <ptredak@nvidia.com>
[ROCm/rccl commit: c05312f151]
Tento commit je obsažen v:
Sylvain Jeaugey
odevzdal
Przemek Tredak
Przemek Tredak
rodič
a867db80f1
revize
42c3edc35a
@@ -1,238 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* Redistribution and use in source and binary forms, with or without
|
||||
* modification, are permitted provided that the following conditions
|
||||
* are met:
|
||||
* * Redistributions of source code must retain the above copyright
|
||||
* notice, this list of conditions and the following disclaimer.
|
||||
* * Redistributions in binary form must reproduce the above copyright
|
||||
* notice, this list of conditions and the following disclaimer in the
|
||||
* documentation and/or other materials provided with the distribution.
|
||||
* * Neither the name of NVIDIA CORPORATION nor the names of its
|
||||
* contributors may be used to endorse or promote products derived
|
||||
* from this software without specific prior written permission.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
|
||||
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
|
||||
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
|
||||
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
|
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
|
||||
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
|
||||
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
************************************************************************/
|
||||
|
||||
#include <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "nccl.h"
|
||||
#include "test_utilities.h"
|
||||
|
||||
|
||||
template<typename T>
|
||||
void RunTest(T** sendbuff, T** recvbuff, const int N, const ncclDataType_t type,
|
||||
ncclComm_t* const comms, const std::vector<int>& dList) {
|
||||
// initialize data
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
|
||||
T* buffer = (T*)malloc(nDev * N * sizeof(T));
|
||||
T* result = (T*)malloc(nDev * N * sizeof(T));
|
||||
memset(buffer, 0, nDev * N * sizeof(T));
|
||||
memset(result, 0, nDev * N * sizeof(T));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamCreate(s+i));
|
||||
CUDACHECK(cudaMemset(recvbuff[i], 0, nDev * N * sizeof(T)));
|
||||
Randomize(sendbuff[i], N, i);
|
||||
|
||||
CUDACHECK(cudaMemcpy(result + i * N, sendbuff[i], N * sizeof(T),
|
||||
cudaMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
// warm up GPU
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclAllGather((const void*)sendbuff[i], std::min(32 * 1024, N), type,
|
||||
(void*)recvbuff[i], comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
//for (int n = 1; n <= N; n = n << 1)
|
||||
{
|
||||
int n = N;
|
||||
printf("%12i %12i %6s", (int)(n * sizeof(T)), n, TypeName(type).c_str());
|
||||
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclAllGather((const void*)sendbuff[i], n, type, (void*)recvbuff[i], comms[i],
|
||||
s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count();
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 * (double)(nDev - 1)
|
||||
/ elapsedSec;
|
||||
double busbw = algbw;
|
||||
|
||||
double maxDelta = 0.0;
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
double tmpDelta = CheckDelta<T>(recvbuff[i], result, nDev*N);
|
||||
maxDelta = std::max(tmpDelta, maxDelta);
|
||||
}
|
||||
|
||||
printf(" %7.3f %5.2f %5.2f %7.0le\n", elapsedSec * 1.0E3, algbw, busbw,
|
||||
maxDelta);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamDestroy(s[i]));
|
||||
}
|
||||
free(s);
|
||||
free(buffer);
|
||||
free(result);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void RunTests(const int N, const ncclDataType_t type, ncclComm_t* const comms,
|
||||
const std::vector<int>& dList) {
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
T** sendbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
T** recvbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaMalloc(sendbuff + i, N * sizeof(T)));
|
||||
CUDACHECK(cudaMalloc(recvbuff + i, nDev * N * sizeof(T)));
|
||||
}
|
||||
|
||||
RunTest<T>(sendbuff, recvbuff, N, type, comms, dList);
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaFree(sendbuff[i]));
|
||||
CUDACHECK(cudaFree(recvbuff[i]));
|
||||
}
|
||||
|
||||
free(sendbuff);
|
||||
free(recvbuff);
|
||||
}
|
||||
|
||||
void usage() {
|
||||
printf("Tests nccl AllGather with user supplied arguments.\n"
|
||||
" Usage: all_reduce_test <data size in bytes> [number of GPUs] "
|
||||
"[GPU 0] [GPU 1] ...\n\n");
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
int nVis = 0;
|
||||
CUDACHECK(cudaGetDeviceCount(&nVis));
|
||||
|
||||
int N = 0;
|
||||
if (argc > 1) {
|
||||
int t = sscanf(argv[1], "%d", &N);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
} else {
|
||||
printf("Error: must specify at least data size in bytes!\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
int nDev = nVis;
|
||||
if (argc > 2) {
|
||||
int t = sscanf(argv[2], "%d", &nDev);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
std::vector<int> dList(nDev);
|
||||
for (int i = 0; i < nDev; ++i)
|
||||
dList[i] = i % nVis;
|
||||
|
||||
|
||||
if (argc > 3) {
|
||||
if (argc - 3 != nDev) {
|
||||
printf("Error: insufficient number of GPUs in list\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
int t = sscanf(argv[3 + i], "%d", dList.data() + i);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[2 + i]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ncclComm_t* comms = (ncclComm_t*)malloc(sizeof(ncclComm_t)*nDev);
|
||||
ncclCommInitAll(comms, nDev, dList.data());
|
||||
|
||||
printf("# Using devices\n");
|
||||
for (int g=0; g<nDev; ++g) {
|
||||
int cudaDev;
|
||||
int rank;
|
||||
cudaDeviceProp prop;
|
||||
ncclCommCuDevice(comms[g], &cudaDev);
|
||||
ncclCommUserRank(comms[g], &rank);
|
||||
CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
|
||||
printf("# Rank %2d uses device %2d [0x%02x] %s\n", rank, cudaDev,
|
||||
prop.pciBusID, prop.name);
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
printf("# %10s %12s %6s %7s %5s %5s %7s\n",
|
||||
"bytes", "N", "type", "time", "algbw", "busbw", "delta");
|
||||
|
||||
RunTests<char>(N / sizeof(char), ncclChar, comms, dList);
|
||||
RunTests<int>(N / sizeof(int), ncclInt, comms, dList);
|
||||
#if CUDART_VERSION >= 7050
|
||||
RunTests<half>(N / sizeof(half), ncclHalf, comms, dList);
|
||||
#endif
|
||||
RunTests<float>(N / sizeof(float), ncclFloat, comms, dList);
|
||||
RunTests<double>(N / sizeof(double), ncclDouble, comms, dList);
|
||||
RunTests<long long>(N / sizeof(long long), ncclInt64, comms, dList);
|
||||
RunTests<unsigned long long>(N / sizeof(unsigned long long), ncclUint64, comms, dList);
|
||||
|
||||
printf("\n");
|
||||
|
||||
for(int i=0; i<nDev; ++i)
|
||||
ncclCommDestroy(comms[i]);
|
||||
free(comms);
|
||||
|
||||
exit(EXIT_SUCCESS);
|
||||
}
|
||||
|
||||
@@ -1,301 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* Redistribution and use in source and binary forms, with or without
|
||||
* modification, are permitted provided that the following conditions
|
||||
* are met:
|
||||
* * Redistributions of source code must retain the above copyright
|
||||
* notice, this list of conditions and the following disclaimer.
|
||||
* * Redistributions in binary form must reproduce the above copyright
|
||||
* notice, this list of conditions and the following disclaimer in the
|
||||
* documentation and/or other materials provided with the distribution.
|
||||
* * Neither the name of NVIDIA CORPORATION nor the names of its
|
||||
* contributors may be used to endorse or promote products derived
|
||||
* from this software without specific prior written permission.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
|
||||
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
|
||||
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
|
||||
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
|
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
|
||||
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
|
||||
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
************************************************************************/
|
||||
|
||||
#include <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "nccl.h"
|
||||
#include "test_utilities.h"
|
||||
#include <nvToolsExt.h>
|
||||
|
||||
int csv = false;
|
||||
|
||||
template<typename T>
|
||||
void RunTest(T** sendbuff, T** recvbuff, const int N, const ncclDataType_t type,
|
||||
const ncclRedOp_t op, ncclComm_t* comms, const std::vector<int>& dList) {
|
||||
// initialize data
|
||||
T* buffer = (T*)malloc(N * sizeof(T));
|
||||
T* result = (T*)malloc(N * sizeof(T));
|
||||
memset(buffer, 0, N * sizeof(T));
|
||||
memset(result, 0, N * sizeof(T));
|
||||
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamCreate(s+i));
|
||||
CUDACHECK(cudaMemset(recvbuff[i], 0, N * sizeof(T)));
|
||||
Randomize(sendbuff[i], N, i);
|
||||
if(i == 0) {
|
||||
CUDACHECK(cudaMemcpy(result, sendbuff[i], N*sizeof(T), cudaMemcpyDeviceToHost));
|
||||
} else {
|
||||
Accumulate<T>(result, sendbuff[i], N, op);
|
||||
}
|
||||
}
|
||||
|
||||
// warm up GPU
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], std::min(N, 1024 * 1024), type, op, comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
// for (int n = 0; n <= N; n = (n > 0) ? n << 1 : 1)
|
||||
{
|
||||
int n = N;
|
||||
printf((csv) ? "%i,%i,%s,%s," : "%12i %12i %6s %6s",
|
||||
(int) (n * sizeof(T)), n, TypeName(type).c_str(),
|
||||
OperationName(op).c_str());
|
||||
|
||||
// do out-of-place reduction first
|
||||
nvtxRangePushA("out of place");
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
//for (int i=0; i<100; i++) {
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], n, type, op,
|
||||
comms[i], s[i]);
|
||||
}
|
||||
//}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
nvtxRangePop();
|
||||
|
||||
nvtxRangePushA("out of place bookkeeping");
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count(); // / 100.0;
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 / elapsedSec;
|
||||
double busbw = algbw * (double)(2 * nDev - 2) / (double)nDev;
|
||||
|
||||
double maxDelta = 0.0;
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
double tmpDelta = CheckDelta<T>(recvbuff[i], result, N);
|
||||
maxDelta = std::max(tmpDelta, maxDelta);
|
||||
}
|
||||
|
||||
printf((csv)?"%f,%f,%f,%le,":" %7.3f %5.2f %5.2f %7.0le",
|
||||
elapsedSec * 1.0E3, algbw, busbw, maxDelta);
|
||||
|
||||
nvtxRangePop();
|
||||
}
|
||||
|
||||
|
||||
// for (int n = 0; n <= N; n = (n > 0) ? n << 1 : 1)
|
||||
{
|
||||
int n = N;
|
||||
// now do in-place reduction
|
||||
nvtxRangePushA("in place");
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
//for (int i=0; i<100; i++) {
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclAllReduce((const void*)sendbuff[i], (void*)sendbuff[i], n, type, op,
|
||||
comms[i], s[i]);
|
||||
}
|
||||
//}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
nvtxRangePop();
|
||||
|
||||
nvtxRangePushA("in place bookkeeping");
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count(); // / 100.0;
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 / elapsedSec;
|
||||
double busbw = algbw * (double)(2 * nDev - 2) / (double)nDev;
|
||||
|
||||
double maxDelta = 0.0;
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
double tmpDelta = CheckDelta<T>(sendbuff[i], result, N);
|
||||
maxDelta = std::max(tmpDelta, maxDelta);
|
||||
}
|
||||
|
||||
printf((csv)?"%f,%f,%f,%le,":" %7.3f %5.2f %5.2f %7.0le\n",
|
||||
elapsedSec * 1.0E3, algbw, busbw, maxDelta);
|
||||
|
||||
nvtxRangePop();
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamDestroy(s[i]));
|
||||
}
|
||||
free(s);
|
||||
free(buffer);
|
||||
free(result);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void RunTests(const int N, const ncclDataType_t type, ncclComm_t* comms,
|
||||
const std::vector<int>& dList) {
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
T** sendbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
T** recvbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaMalloc(sendbuff + i, N * sizeof(T)));
|
||||
CUDACHECK(cudaMalloc(recvbuff + i, N * sizeof(T)));
|
||||
}
|
||||
|
||||
for (ncclRedOp_t op : { ncclSum, ncclProd, ncclMax, ncclMin }) {
|
||||
// for (ncclRedOp_t op : { ncclSum }) {
|
||||
RunTest<T>(sendbuff, recvbuff, N, type, op, comms, dList);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaFree(sendbuff[i]));
|
||||
CUDACHECK(cudaFree(recvbuff[i]));
|
||||
}
|
||||
|
||||
free(sendbuff);
|
||||
free(recvbuff);
|
||||
}
|
||||
|
||||
void usage() {
|
||||
printf("Tests nccl AllReduce with user supplied arguments.\n"
|
||||
" Usage: all_reduce_test <data size in bytes> [number of GPUs] "
|
||||
"[GPU 0] [GPU 1] ...\n\n");
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
int nVis = 0;
|
||||
CUDACHECK(cudaGetDeviceCount(&nVis));
|
||||
|
||||
int N = 0;
|
||||
if (argc > 1) {
|
||||
int t = sscanf(argv[1], "%d", &N);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
} else {
|
||||
printf("Error: must specify at least data size in bytes!\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
int nDev = nVis;
|
||||
if (argc > 2) {
|
||||
int t = sscanf(argv[2], "%d", &nDev);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
std::vector<int> dList(nDev);
|
||||
for (int i = 0; i < nDev; ++i)
|
||||
dList[i] = i % nVis;
|
||||
|
||||
if (argc > 3) {
|
||||
if (argc - 3 != nDev) {
|
||||
printf("Error: insufficient number of GPUs in list\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
int t = sscanf(argv[3 + i], "%d", dList.data() + i);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[2 + i]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ncclComm_t* comms = (ncclComm_t*)malloc(sizeof(ncclComm_t)*nDev);
|
||||
ncclCommInitAll(comms, nDev, dList.data());
|
||||
|
||||
if (!csv) {
|
||||
printf("# Using devices\n");
|
||||
for (int g = 0; g < nDev; ++g) {
|
||||
int cudaDev;
|
||||
int rank;
|
||||
cudaDeviceProp prop;
|
||||
ncclCommCuDevice(comms[g], &cudaDev);
|
||||
ncclCommUserRank(comms[g], &rank);
|
||||
CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
|
||||
printf("# Rank %2d uses device %2d [0x%02x] %s\n", rank, cudaDev,
|
||||
prop.pciBusID, prop.name);
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
printf("# %10s %12s %6s %6s out-of-place in-place\n", "", "", "", "");
|
||||
printf("# %10s %12s %6s %6s %7s %5s %5s %7s %7s %5s %5s %7s\n", "bytes", "N", "type", "op",
|
||||
"time", "algbw", "busbw", "res", "time", "algbw", "busbw", "res");
|
||||
}
|
||||
else {
|
||||
printf("B,N,type,op,oop_time,oop_algbw,oop_busbw,oop_res,ip_time,ip_algbw,ip_busbw,ip_res\n");
|
||||
}
|
||||
|
||||
RunTests<char>(N / sizeof(char), ncclChar, comms, dList);
|
||||
RunTests<int>(N / sizeof(int), ncclInt, comms, dList);
|
||||
#if CUDART_VERSION >= 7050
|
||||
RunTests<half>(N / sizeof(half), ncclHalf, comms, dList);
|
||||
#endif
|
||||
RunTests<float>(N / sizeof(float), ncclFloat, comms, dList);
|
||||
RunTests<double>(N / sizeof(double), ncclDouble, comms, dList);
|
||||
RunTests<long long>(N / sizeof(long long), ncclInt64, comms, dList);
|
||||
RunTests<unsigned long long>(N / sizeof(unsigned long long), ncclUint64, comms, dList);
|
||||
|
||||
printf("\n");
|
||||
|
||||
for(int i=0; i<nDev; ++i)
|
||||
ncclCommDestroy(comms[i]);
|
||||
free(comms);
|
||||
|
||||
exit(EXIT_SUCCESS);
|
||||
}
|
||||
|
||||
@@ -1,238 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* Redistribution and use in source and binary forms, with or without
|
||||
* modification, are permitted provided that the following conditions
|
||||
* are met:
|
||||
* * Redistributions of source code must retain the above copyright
|
||||
* notice, this list of conditions and the following disclaimer.
|
||||
* * Redistributions in binary form must reproduce the above copyright
|
||||
* notice, this list of conditions and the following disclaimer in the
|
||||
* documentation and/or other materials provided with the distribution.
|
||||
* * Neither the name of NVIDIA CORPORATION nor the names of its
|
||||
* contributors may be used to endorse or promote products derived
|
||||
* from this software without specific prior written permission.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
|
||||
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
|
||||
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
|
||||
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
|
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
|
||||
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
|
||||
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
************************************************************************/
|
||||
|
||||
#include <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "nccl.h"
|
||||
#include "test_utilities.h"
|
||||
|
||||
|
||||
template<typename T>
|
||||
void RunTest(T** buff, const int N, const ncclDataType_t type, const int root,
|
||||
ncclComm_t* const comms, const std::vector<int>& dList) {
|
||||
// initialize data
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
|
||||
T* buffer = (T*)malloc(N * sizeof(T));
|
||||
T* result = (T*)malloc(N * sizeof(T));
|
||||
memset(result, 0, N * sizeof(T));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamCreate(s+i));
|
||||
|
||||
if (i == root) {
|
||||
Randomize(buff[root], N, root);
|
||||
CUDACHECK(cudaMemcpy(result, buff[root], N * sizeof(T),
|
||||
cudaMemcpyDeviceToHost));
|
||||
} else {
|
||||
CUDACHECK(cudaMemset(buff[i], 0, N * sizeof(T)));
|
||||
}
|
||||
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
|
||||
// warm up GPU
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclBcast((void*)buff[i], std::min(32 * 1024, N), type, root, comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
// for (int n = 1; n <= N; n = n << 1)
|
||||
{
|
||||
int n = N;
|
||||
printf("%12i %12i %6s %4i", (int)(n * sizeof(T)), n,
|
||||
TypeName(type).c_str(), root);
|
||||
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclBcast((void*)buff[i], n, type, root, comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count();
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 / elapsedSec;
|
||||
double busbw = algbw;
|
||||
|
||||
double maxDelta = 0.0;
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
double tmpDelta = CheckDelta<T>(buff[i], result, n);
|
||||
maxDelta = std::max(tmpDelta, maxDelta);
|
||||
}
|
||||
|
||||
printf(" %7.3f %5.2f %5.2f %7.0le\n", elapsedSec * 1.0E3, algbw, busbw,
|
||||
maxDelta);
|
||||
}
|
||||
|
||||
for(int i=0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamDestroy(s[i]));
|
||||
}
|
||||
free(s);
|
||||
free(buffer);
|
||||
free(result);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void RunTests(const int N, const ncclDataType_t type, ncclComm_t* const comms,
|
||||
const std::vector<int>& dList) {
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
T** buff = (T**)malloc(nDev * sizeof(T*));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaMalloc(buff + i, N * sizeof(T)));
|
||||
}
|
||||
|
||||
//for (int root = 1; root < 2; ++root) {
|
||||
for (int root = 0; root < nDev; ++root) {
|
||||
RunTest<T>(buff, N, type, root, comms, dList);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaFree(buff[i]));
|
||||
}
|
||||
|
||||
free(buff);
|
||||
}
|
||||
|
||||
void usage() {
|
||||
printf("Tests nccl Broadcast with user supplied arguments.\n"
|
||||
" Usage: broadcast_test <data size in bytes> [number of GPUs] "
|
||||
"[GPU 0] [GPU 1] ...\n\n");
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
int nVis = 0;
|
||||
CUDACHECK(cudaGetDeviceCount(&nVis));
|
||||
|
||||
unsigned long long N = 0;
|
||||
if (argc > 1) {
|
||||
int t = sscanf(argv[1], "%llu", &N);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
} else {
|
||||
printf("Error: must specify at least data size in bytes!\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
int nDev = nVis;
|
||||
if (argc > 2) {
|
||||
int t = sscanf(argv[2], "%d", &nDev);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
std::vector<int> dList(nDev);
|
||||
for (int i = 0; i < nDev; ++i)
|
||||
dList[i] = i % nVis;
|
||||
|
||||
if (argc > 3) {
|
||||
if (argc - 3 != nDev) {
|
||||
printf("Error: insufficient number of GPUs in list\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
int t = sscanf(argv[3 + i], "%d", dList.data() + i);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[2 + i]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ncclComm_t* comms = (ncclComm_t*)malloc(sizeof(ncclComm_t)*nDev);;
|
||||
ncclCommInitAll(comms, nDev, dList.data());
|
||||
|
||||
printf("# Using devices\n");
|
||||
for (int g = 0; g < nDev; ++g) {
|
||||
int cudaDev;
|
||||
int rank;
|
||||
cudaDeviceProp prop;
|
||||
ncclCommCuDevice(comms[g], &cudaDev);
|
||||
ncclCommUserRank(comms[g], &rank);
|
||||
CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
|
||||
printf("# Rank %2d uses device %2d [0x%02x] %s\n", rank, cudaDev,
|
||||
prop.pciBusID, prop.name);
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
printf("# %10s %12s %6s %4s %7s %5s %5s %7s\n",
|
||||
"bytes", "N", "type", "root", "time", "algbw", "busbw", "delta");
|
||||
|
||||
RunTests<char>(N / sizeof(char), ncclChar, comms, dList);
|
||||
RunTests<int>(N / sizeof(int), ncclInt, comms, dList);
|
||||
#ifdef CUDA_HAS_HALF
|
||||
RunTests<half>(N / sizeof(half), ncclHalf, comms, dList);
|
||||
#endif
|
||||
RunTests<float>(N / sizeof(float), ncclFloat, comms, dList);
|
||||
RunTests<double>(N / sizeof(double), ncclDouble, comms, dList);
|
||||
RunTests<long long>(N / sizeof(long long), ncclInt64, comms, dList);
|
||||
RunTests<unsigned long long>(N / sizeof(unsigned long long), ncclUint64, comms, dList);
|
||||
|
||||
printf("\n");
|
||||
|
||||
for(int i = 0; i < nDev; ++i)
|
||||
ncclCommDestroy(comms[i]);
|
||||
free(comms);
|
||||
|
||||
exit(EXIT_SUCCESS);
|
||||
}
|
||||
|
||||
@@ -1,83 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* Redistribution and use in source and binary forms, with or without
|
||||
* modification, are permitted provided that the following conditions
|
||||
* are met:
|
||||
* * Redistributions of source code must retain the above copyright
|
||||
* notice, this list of conditions and the following disclaimer.
|
||||
* * Redistributions in binary form must reproduce the above copyright
|
||||
* notice, this list of conditions and the following disclaimer in the
|
||||
* documentation and/or other materials provided with the distribution.
|
||||
* * Neither the name of NVIDIA CORPORATION nor the names of its
|
||||
* contributors may be used to endorse or promote products derived
|
||||
* from this software without specific prior written permission.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
|
||||
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
|
||||
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
|
||||
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
|
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
|
||||
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
|
||||
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
************************************************************************/
|
||||
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
|
||||
#include "nccl.h"
|
||||
#include "mpi.h"
|
||||
|
||||
#define CUDACHECK(cmd) do { \
|
||||
cudaError_t e = cmd; \
|
||||
if( e != cudaSuccess ) { \
|
||||
printf("Cuda failure %s:%d '%s'\n", \
|
||||
__FILE__,__LINE__,cudaGetErrorString(e)); \
|
||||
exit(EXIT_FAILURE); \
|
||||
} \
|
||||
} while(false)
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
ncclUniqueId commId;
|
||||
int size, rank;
|
||||
int ret;
|
||||
|
||||
MPI_Init(&argc, &argv);
|
||||
MPI_Comm_size(MPI_COMM_WORLD, &size);
|
||||
MPI_Comm_rank(MPI_COMM_WORLD, &rank);
|
||||
|
||||
int gpu = atoi(argv[rank+1]);
|
||||
printf("MPI Rank %d running on GPU %d\n", rank, gpu);
|
||||
// We have to set our device before NCCL init
|
||||
CUDACHECK(cudaSetDevice(gpu));
|
||||
MPI_Barrier(MPI_COMM_WORLD);
|
||||
|
||||
ncclComm_t comm;
|
||||
// Let's use rank 0 PID as job ID
|
||||
ncclGetUniqueId(&commId);
|
||||
MPI_Bcast(&commId, NCCL_UNIQUE_ID_BYTES, MPI_CHAR, 0, MPI_COMM_WORLD);
|
||||
ret = ncclCommInitRank(&comm, size, commId, rank);
|
||||
if (ret != ncclSuccess) {
|
||||
printf("NCCL Init failed : %d\n", ret);
|
||||
exit(1);
|
||||
}
|
||||
|
||||
int *dptr;
|
||||
CUDACHECK(cudaMalloc(&dptr, 1024*2*sizeof(int)));
|
||||
int val = rank;
|
||||
CUDACHECK(cudaMemcpy(dptr, &val, sizeof(int), cudaMemcpyHostToDevice));
|
||||
|
||||
ncclAllReduce((const void*)dptr, (void*)(dptr+1024), 1024, ncclInt, ncclSum, comm, cudaStreamDefault);
|
||||
|
||||
CUDACHECK(cudaMemcpy(&val, (dptr+1024), sizeof(int), cudaMemcpyDeviceToHost));
|
||||
printf("Sum is %d\n", val);
|
||||
CUDACHECK(cudaFree(dptr));
|
||||
|
||||
MPI_Finalize();
|
||||
ncclCommDestroy(comm);
|
||||
return 0;
|
||||
}
|
||||
@@ -1,285 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* Redistribution and use in source and binary forms, with or without
|
||||
* modification, are permitted provided that the following conditions
|
||||
* are met:
|
||||
* * Redistributions of source code must retain the above copyright
|
||||
* notice, this list of conditions and the following disclaimer.
|
||||
* * Redistributions in binary form must reproduce the above copyright
|
||||
* notice, this list of conditions and the following disclaimer in the
|
||||
* documentation and/or other materials provided with the distribution.
|
||||
* * Neither the name of NVIDIA CORPORATION nor the names of its
|
||||
* contributors may be used to endorse or promote products derived
|
||||
* from this software without specific prior written permission.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
|
||||
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
|
||||
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
|
||||
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
|
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
|
||||
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
|
||||
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
************************************************************************/
|
||||
|
||||
#include <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "nccl.h"
|
||||
#include "test_utilities.h"
|
||||
|
||||
|
||||
template<typename T>
|
||||
void RunTest(T** sendbuff, T** recvbuff, const int N, const ncclDataType_t type,
|
||||
const ncclRedOp_t op, ncclComm_t* const comms, const std::vector<int>& dList) {
|
||||
// initialize data
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
|
||||
|
||||
T* buffer = (T*)malloc(N * nDev * sizeof(T));
|
||||
T* result = (T*)malloc(N * nDev * sizeof(T));
|
||||
memset(buffer, 0, N * nDev * sizeof(T));
|
||||
memset(result, 0, N * nDev * sizeof(T));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamCreate(s+i));
|
||||
CUDACHECK(cudaMemset(recvbuff[i], 0, N * sizeof(T)));
|
||||
Randomize(sendbuff[i], N * nDev, i);
|
||||
|
||||
if (i == 0) {
|
||||
CUDACHECK(cudaMemcpy(result, sendbuff[i], N * nDev * sizeof(T),
|
||||
cudaMemcpyDeviceToHost));
|
||||
} else {
|
||||
Accumulate<T>(result, sendbuff[i], N * nDev, op);
|
||||
}
|
||||
}
|
||||
|
||||
// warm up GPU
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclReduceScatter((const void*)sendbuff[i], (void*)recvbuff[i],
|
||||
std::min(N, 1024 * 1024), type, op, comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
// for (int n = 0; n <= N; n = (n > 0) ? n << 1 : 1)
|
||||
{
|
||||
int n = N;
|
||||
printf("%12i %12i %6s %6s", (int)(n * sizeof(T)), n,
|
||||
TypeName(type).c_str(), OperationName(op).c_str());
|
||||
|
||||
// do out-of-place reduction first
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclReduceScatter((const void*)sendbuff[i], (void*)recvbuff[i], n, type,
|
||||
op, comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count();
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 / elapsedSec;
|
||||
double busbw = algbw * (double)(nDev - 1);
|
||||
|
||||
double maxDelta = 0.0;
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
double tmpDelta = CheckDelta<T>(recvbuff[i], result+i*n, n);
|
||||
maxDelta = std::max(tmpDelta, maxDelta);
|
||||
}
|
||||
|
||||
printf(" %7.3f %5.2f %5.2f %7.0le", elapsedSec * 1.0E3, algbw, busbw,
|
||||
maxDelta);
|
||||
}
|
||||
|
||||
{
|
||||
// now do in-place reduction
|
||||
int n = N;
|
||||
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclReduceScatter((const void*)sendbuff[i], (void*)sendbuff[i], n, type,
|
||||
op, comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count();
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 / elapsedSec;
|
||||
double busbw = algbw * (double)(nDev - 1);
|
||||
|
||||
double maxDelta = 0.0;
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
double tmpDelta = CheckDelta<T>(sendbuff[i], result+i*n, n);
|
||||
maxDelta = std::max(tmpDelta, maxDelta);
|
||||
}
|
||||
|
||||
printf(" %7.3f %5.2f %5.2f %7.0le\n", elapsedSec * 1.0E3, algbw, busbw,
|
||||
maxDelta);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamDestroy(s[i]));
|
||||
}
|
||||
free(s);
|
||||
free(buffer);
|
||||
free(result);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void RunTests(const int N, const ncclDataType_t type, ncclComm_t* const comms,
|
||||
const std::vector<int>& dList) {
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
T** sendbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
T** recvbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaMalloc(sendbuff + i, N * nDev * sizeof(T)));
|
||||
CUDACHECK(cudaMalloc(recvbuff + i, N * sizeof(T)));
|
||||
}
|
||||
|
||||
for (ncclRedOp_t op : { ncclSum, ncclProd, ncclMax, ncclMin }) {
|
||||
// for (ncclRedOp_t op : { ncclSum }) {
|
||||
RunTest<T>(sendbuff, recvbuff, N, type, op, comms, dList);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaFree(sendbuff[i]));
|
||||
CUDACHECK(cudaFree(recvbuff[i]));
|
||||
}
|
||||
|
||||
free(sendbuff);
|
||||
free(recvbuff);
|
||||
}
|
||||
|
||||
void usage() {
|
||||
printf("Tests nccl ReduceScatter with user supplied arguments.\n"
|
||||
" Usage: all_reduce_test <data size in bytes> [number of GPUs] "
|
||||
"[GPU 0] [GPU 1] ...\n\n");
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
int nVis = 0;
|
||||
CUDACHECK(cudaGetDeviceCount(&nVis));
|
||||
|
||||
int N = 0;
|
||||
if (argc > 1) {
|
||||
int t = sscanf(argv[1], "%d", &N);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
} else {
|
||||
printf("Error: must specify at least data size in bytes!\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
int nDev = nVis;
|
||||
if (argc > 2) {
|
||||
int t = sscanf(argv[2], "%d", &nDev);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
std::vector<int> dList(nDev);
|
||||
for (int i = 0; i < nDev; ++i)
|
||||
dList[i] = i % nVis;
|
||||
|
||||
if (argc > 3) {
|
||||
if (argc - 3 != nDev) {
|
||||
printf("Error: insufficient number of GPUs in list\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
int t = sscanf(argv[3 + i], "%d", dList.data() + i);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[2 + i]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ncclComm_t* comms = (ncclComm_t*)malloc(sizeof(ncclComm_t)*nDev);
|
||||
ncclCommInitAll(comms, nDev, dList.data());
|
||||
|
||||
printf("# Using devices\n");
|
||||
for (int g = 0; g < nDev; ++g) {
|
||||
int cudaDev;
|
||||
int rank;
|
||||
cudaDeviceProp prop;
|
||||
ncclCommCuDevice(comms[g], &cudaDev);
|
||||
ncclCommUserRank(comms[g], &rank);
|
||||
CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
|
||||
printf("# Rank %2d uses device %2d [0x%02x] %s\n", rank, cudaDev,
|
||||
prop.pciBusID, prop.name);
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
printf("# %10s %12s %6s %6s out-of-place "
|
||||
"in-place\n", "", "", "", "");
|
||||
printf("# %10s %12s %6s %6s %7s %5s %5s %7s %7s %5s %5s %7s\n",
|
||||
"bytes", "N", "type", "op", "time", "algbw", "busbw", "delta", "time",
|
||||
"algbw", "busbw", "delta");
|
||||
|
||||
RunTests<char>(N / sizeof(char), ncclChar, comms, dList);
|
||||
RunTests<int>(N / sizeof(int), ncclInt, comms, dList);
|
||||
#ifdef CUDA_HAS_HALF
|
||||
RunTests<half>(N / sizeof(half), ncclHalf, comms, dList);
|
||||
#endif
|
||||
RunTests<float>(N / sizeof(float), ncclFloat, comms, dList);
|
||||
RunTests<double>(N / sizeof(double), ncclDouble, comms, dList);
|
||||
RunTests<long long>(N / sizeof(long long), ncclInt64, comms, dList);
|
||||
RunTests<unsigned long long>(N / sizeof(unsigned long long), ncclUint64, comms, dList);
|
||||
|
||||
printf("\n");
|
||||
|
||||
for(int i=0; i<nDev; ++i)
|
||||
ncclCommDestroy(comms[i]);
|
||||
free(comms);
|
||||
|
||||
exit(EXIT_SUCCESS);
|
||||
}
|
||||
|
||||
@@ -1,299 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* Redistribution and use in source and binary forms, with or without
|
||||
* modification, are permitted provided that the following conditions
|
||||
* are met:
|
||||
* * Redistributions of source code must retain the above copyright
|
||||
* notice, this list of conditions and the following disclaimer.
|
||||
* * Redistributions in binary form must reproduce the above copyright
|
||||
* notice, this list of conditions and the following disclaimer in the
|
||||
* documentation and/or other materials provided with the distribution.
|
||||
* * Neither the name of NVIDIA CORPORATION nor the names of its
|
||||
* contributors may be used to endorse or promote products derived
|
||||
* from this software without specific prior written permission.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
|
||||
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
|
||||
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
|
||||
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
|
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
|
||||
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
|
||||
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
************************************************************************/
|
||||
|
||||
#include <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "nccl.h"
|
||||
#include "test_utilities.h"
|
||||
#include <nvToolsExt.h>
|
||||
|
||||
int csv = false;
|
||||
|
||||
template<typename T>
|
||||
void RunTest(T** sendbuff, T** recvbuff, const int N, const ncclDataType_t type,
|
||||
const ncclRedOp_t op, int root, ncclComm_t* const comms,
|
||||
const std::vector<int>& dList) {
|
||||
|
||||
// initialize data
|
||||
T* buffer = (T*)malloc(N * sizeof(T));
|
||||
T* result = (T*)malloc(N * sizeof(T));
|
||||
memset(buffer, 0, N * sizeof(T));
|
||||
memset(result, 0, N * sizeof(T));
|
||||
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamCreate(s+i));
|
||||
CUDACHECK(cudaMemset(recvbuff[i], 0, N * sizeof(T)));
|
||||
Randomize(sendbuff[i], N, i);
|
||||
if(i == 0) {
|
||||
CUDACHECK(cudaMemcpy(result, sendbuff[i], N*sizeof(T), cudaMemcpyDeviceToHost));
|
||||
} else {
|
||||
Accumulate<T>(result, sendbuff[i], N, op);
|
||||
}
|
||||
}
|
||||
|
||||
// warm up GPU
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclReduce((const void*)sendbuff[i], (void*)recvbuff[i], std::min(N, 1024 * 1024),
|
||||
type, op, root, comms[i], s[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
// for (int n = 0; n <= N; n = (n > 0) ? n << 1 : 1)
|
||||
{
|
||||
int n = N;
|
||||
printf((csv) ? "%i,%i,%s,%s,%d," : "%12i %12i %6s %6s %4d",
|
||||
(int) (n * sizeof(T)), n, TypeName(type).c_str(),
|
||||
OperationName(op).c_str(), root);
|
||||
|
||||
// do out-of-place reduction first
|
||||
nvtxRangePushA("out of place");
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
//for (int i=0; i<100; i++) {
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclReduce((const void*)sendbuff[i], (void*)recvbuff[i], n, type, op,
|
||||
root, comms[i], s[i]);
|
||||
}
|
||||
//}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
nvtxRangePop();
|
||||
|
||||
nvtxRangePushA("out of place bookkeeping");
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count(); // / 100.0;
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 / elapsedSec;
|
||||
double busbw = algbw;
|
||||
|
||||
CUDACHECK(cudaSetDevice(dList[root]));
|
||||
double maxDelta = CheckDelta<T>(recvbuff[root], result, N);
|
||||
|
||||
printf((csv)?"%f,%f,%f,%le,":" %7.3f %5.2f %5.2f %7.0le",
|
||||
elapsedSec * 1.0E3, algbw, busbw, maxDelta);
|
||||
|
||||
nvtxRangePop();
|
||||
}
|
||||
|
||||
|
||||
// for (int n = 0; n <= N; n = (n > 0) ? n << 1 : 1)
|
||||
{
|
||||
int n = N;
|
||||
// now do in-place reduction
|
||||
nvtxRangePushA("in place");
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
//for (int i=0; i<100; i++) {
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
ncclReduce((const void*)sendbuff[i], (void*)sendbuff[i], n, type, op,
|
||||
root, comms[i], s[i]);
|
||||
}
|
||||
//}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamSynchronize(s[i]));
|
||||
}
|
||||
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
nvtxRangePop();
|
||||
|
||||
nvtxRangePushA("in place bookkeeping");
|
||||
double elapsedSec =
|
||||
std::chrono::duration_cast<std::chrono::duration<double>>(
|
||||
stop - start).count(); // / 100.0;
|
||||
double algbw = (double)(n * sizeof(T)) / 1.0E9 / elapsedSec;
|
||||
double busbw = algbw;
|
||||
|
||||
CUDACHECK(cudaSetDevice(dList[root]));
|
||||
double maxDelta = CheckDelta<T>(sendbuff[root], result, N);
|
||||
|
||||
printf((csv)?"%f,%f,%f,%le,":" %7.3f %5.2f %5.2f %7.0le\n",
|
||||
elapsedSec * 1.0E3, algbw, busbw, maxDelta);
|
||||
|
||||
nvtxRangePop();
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaStreamDestroy(s[i]));
|
||||
}
|
||||
free(s);
|
||||
free(buffer);
|
||||
free(result);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void RunTests(const int N, const ncclDataType_t type, ncclComm_t* const comms,
|
||||
const std::vector<int>& dList) {
|
||||
int nDev = 0;
|
||||
ncclCommCount(comms[0], &nDev);
|
||||
T** sendbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
T** recvbuff = (T**)malloc(nDev * sizeof(T*));
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaMalloc(sendbuff + i, N * sizeof(T)));
|
||||
CUDACHECK(cudaMalloc(recvbuff + i, N * sizeof(T)));
|
||||
}
|
||||
|
||||
for (ncclRedOp_t op : { ncclSum, ncclProd, ncclMax, ncclMin }) {
|
||||
// for (ncclRedOp_t op : { ncclSum }) {
|
||||
for(int root=0; root<nDev; ++root) {
|
||||
RunTest<T>(sendbuff, recvbuff, N, type, op, root, comms, dList);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
CUDACHECK(cudaSetDevice(dList[i]));
|
||||
CUDACHECK(cudaFree(sendbuff[i]));
|
||||
CUDACHECK(cudaFree(recvbuff[i]));
|
||||
}
|
||||
|
||||
free(sendbuff);
|
||||
free(recvbuff);
|
||||
}
|
||||
|
||||
void usage() {
|
||||
printf("Tests nccl Reduce with user supplied arguments.\n"
|
||||
" Usage: reduce_test <data size in bytes> [number of GPUs] "
|
||||
"[GPU 0] [GPU 1] ...\n\n");
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
int nVis = 0;
|
||||
CUDACHECK(cudaGetDeviceCount(&nVis));
|
||||
|
||||
int N = 0;
|
||||
if (argc > 1) {
|
||||
int t = sscanf(argv[1], "%d", &N);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
} else {
|
||||
printf("Error: must specify at least data size in bytes!\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
int nDev = nVis;
|
||||
if (argc > 2) {
|
||||
int t = sscanf(argv[2], "%d", &nDev);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[1]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
std::vector<int> dList(nDev);
|
||||
for (int i = 0; i < nDev; ++i)
|
||||
dList[i] = i % nVis;
|
||||
|
||||
if (argc > 3) {
|
||||
if (argc - 3 != nDev) {
|
||||
printf("Error: insufficient number of GPUs in list\n\n");
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nDev; ++i) {
|
||||
int t = sscanf(argv[3 + i], "%d", dList.data() + i);
|
||||
if (t == 0) {
|
||||
printf("Error: %s is not an integer!\n\n", argv[2 + i]);
|
||||
usage();
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ncclComm_t* comms = (ncclComm_t*)malloc(sizeof(ncclComm_t)*nDev);
|
||||
ncclCommInitAll(comms, nDev, dList.data());
|
||||
|
||||
if (!csv) {
|
||||
printf("# Using devices\n");
|
||||
for (int g = 0; g < nDev; ++g) {
|
||||
int cudaDev;
|
||||
int rank;
|
||||
cudaDeviceProp prop;
|
||||
ncclCommCuDevice(comms[g], &cudaDev);
|
||||
ncclCommUserRank(comms[g], &rank);
|
||||
CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
|
||||
printf("# Rank %2d uses device %2d [0x%02x] %s\n", rank, cudaDev,
|
||||
prop.pciBusID, prop.name);
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
printf("# %10s %12s %6s %6s %4s out-of-place in-place\n", "", "", "", "", "");
|
||||
printf("# %10s %12s %6s %6s %4s %7s %5s %5s %7s %7s %5s %5s %7s\n",
|
||||
"bytes", "N", "type", "op", "root",
|
||||
"time", "algbw", "busbw", "res", "time", "algbw", "busbw", "res");
|
||||
}
|
||||
else {
|
||||
printf("B,N,type,op,root,oop_time,oop_algbw,oop_busbw,oop_res,ip_time,ip_algbw,ip_busbw,ip_res\n");
|
||||
}
|
||||
|
||||
RunTests<char>(N / sizeof(char), ncclChar, comms, dList);
|
||||
RunTests<int>(N / sizeof(int), ncclInt, comms, dList);
|
||||
#ifdef CUDA_HAS_HALF
|
||||
RunTests<half>(N / sizeof(half), ncclHalf, comms, dList);
|
||||
#endif
|
||||
RunTests<float>(N / sizeof(float), ncclFloat, comms, dList);
|
||||
RunTests<double>(N / sizeof(double), ncclDouble, comms, dList);
|
||||
RunTests<long long>(N / sizeof(long long), ncclInt64, comms, dList);
|
||||
RunTests<unsigned long long>(N / sizeof(unsigned long long), ncclUint64, comms, dList);
|
||||
|
||||
printf("\n");
|
||||
|
||||
for(int i = 0; i < nDev; ++i)
|
||||
ncclCommDestroy(comms[i]);
|
||||
free(comms);
|
||||
|
||||
exit(EXIT_SUCCESS);
|
||||
}
|
||||
|
||||
@@ -1,312 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved.
|
||||
*
|
||||
* Redistribution and use in source and binary forms, with or without
|
||||
* modification, are permitted provided that the following conditions
|
||||
* are met:
|
||||
* * Redistributions of source code must retain the above copyright
|
||||
* notice, this list of conditions and the following disclaimer.
|
||||
* * Redistributions in binary form must reproduce the above copyright
|
||||
* notice, this list of conditions and the following disclaimer in the
|
||||
* documentation and/or other materials provided with the distribution.
|
||||
* * Neither the name of NVIDIA CORPORATION nor the names of its
|
||||
* contributors may be used to endorse or promote products derived
|
||||
* from this software without specific prior written permission.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
|
||||
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
|
||||
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
|
||||
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
|
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
|
||||
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
|
||||
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
************************************************************************/
|
||||
|
||||
|
||||
#ifndef SRC_TEST_UTILITIES_H_
|
||||
#define SRC_TEST_UTILITIES_H_
|
||||
|
||||
#include <curand.h>
|
||||
|
||||
#define CUDACHECK(cmd) do { \
|
||||
cudaError_t e = cmd; \
|
||||
if( e != cudaSuccess ) { \
|
||||
printf("Cuda failure %s:%d '%s'\n", \
|
||||
__FILE__,__LINE__,cudaGetErrorString(e)); \
|
||||
exit(EXIT_FAILURE); \
|
||||
} \
|
||||
} while(false)
|
||||
|
||||
template<typename T>
|
||||
void Randomize(T* const dest, const int N, const int randomSeed);
|
||||
|
||||
template<typename T>
|
||||
void Accumulate(T* dest, const T* contrib, int N, ncclRedOp_t op);
|
||||
|
||||
template<typename T>
|
||||
double CheckDelta(const T* results, const T* expected, int N);
|
||||
|
||||
#define CURAND_CHK(cmd) \
|
||||
do { \
|
||||
curandStatus_t error = (cmd); \
|
||||
if (error != CURAND_STATUS_SUCCESS) { \
|
||||
printf("CuRAND error %i at %s:%i\n", error, __FILE__ , __LINE__); \
|
||||
exit(EXIT_FAILURE); \
|
||||
} \
|
||||
} while (false)
|
||||
|
||||
|
||||
template<typename T>
|
||||
void GenerateRandom(curandGenerator_t generator, T * const dest,
|
||||
const int N);
|
||||
|
||||
template<>
|
||||
void GenerateRandom<char>(curandGenerator_t generator, char * const dest,
|
||||
const int N) {
|
||||
CURAND_CHK(curandGenerate(generator, (unsigned int*)dest,
|
||||
N * sizeof(char) / sizeof(int)));
|
||||
}
|
||||
|
||||
template<>
|
||||
void GenerateRandom<int>(curandGenerator_t generator, int * const dest,
|
||||
const int N) {
|
||||
CURAND_CHK(curandGenerate(generator, (unsigned int*)dest, N));
|
||||
}
|
||||
|
||||
template<>
|
||||
void GenerateRandom<float>(curandGenerator_t generator, float * const dest,
|
||||
const int N) {
|
||||
CURAND_CHK(curandGenerateUniform(generator, dest, N));
|
||||
}
|
||||
|
||||
template<>
|
||||
void GenerateRandom<double>(curandGenerator_t generator, double * const dest,
|
||||
const int N) {
|
||||
CURAND_CHK(curandGenerateUniformDouble(generator, dest, N));
|
||||
}
|
||||
|
||||
template<>
|
||||
void GenerateRandom<unsigned long long>(curandGenerator_t generator, unsigned long long * const dest,
|
||||
const int N) {
|
||||
CURAND_CHK(curandGenerateLongLong(generator, dest, N));
|
||||
}
|
||||
|
||||
|
||||
template<typename T>
|
||||
void Randomize(T* const dest, const int N, const int randomSeed) {
|
||||
curandGenerator_t gen;
|
||||
CURAND_CHK(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MT19937));
|
||||
CURAND_CHK(curandSetPseudoRandomGeneratorSeed(gen, randomSeed));
|
||||
GenerateRandom<T>(gen, dest, N);
|
||||
CURAND_CHK(curandDestroyGenerator(gen));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
|
||||
template<>
|
||||
void Randomize(unsigned long long* const dest, const int N, const int randomSeed) {
|
||||
curandGenerator_t gen;
|
||||
CURAND_CHK(curandCreateGenerator(&gen, CURAND_RNG_QUASI_SOBOL64));
|
||||
GenerateRandom<unsigned long long>(gen, dest, N);
|
||||
CURAND_CHK(curandDestroyGenerator(gen));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
|
||||
template<>
|
||||
void Randomize(long long* const dest, const int N, const int randomSeed) {
|
||||
curandGenerator_t gen;
|
||||
CURAND_CHK(curandCreateGenerator(&gen, CURAND_RNG_QUASI_SOBOL64));
|
||||
GenerateRandom<unsigned long long>(gen, (unsigned long long *)dest, N);
|
||||
CURAND_CHK(curandDestroyGenerator(gen));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
|
||||
#ifdef CUDA_HAS_HALF
|
||||
__global__ void halve(const float * src, half* dest, int N) {
|
||||
for(int tid = threadIdx.x + blockIdx.x*blockDim.x;
|
||||
tid < N; tid += blockDim.x * gridDim.x)
|
||||
dest[tid] = __float2half(src[tid]);
|
||||
}
|
||||
|
||||
template<>
|
||||
void Randomize<half>(half* const dest, const int N, const int randomSeed) {
|
||||
curandGenerator_t gen;
|
||||
CURAND_CHK(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MT19937));
|
||||
CURAND_CHK(curandSetPseudoRandomGeneratorSeed(gen, randomSeed));
|
||||
|
||||
float* temp;
|
||||
CUDACHECK(cudaMalloc(&temp, N*sizeof(float)));
|
||||
GenerateRandom<float>(gen, temp, N);
|
||||
halve<<<128, 512>>>(temp, dest, N);
|
||||
CURAND_CHK(curandDestroyGenerator(gen));
|
||||
CUDACHECK(cudaFree(temp));
|
||||
CUDACHECK(cudaDeviceSynchronize());
|
||||
}
|
||||
#endif
|
||||
|
||||
template<typename T, int OP> __global__ static
|
||||
void accumKern(T* acum, const T* contrib, int N) {
|
||||
int tid = threadIdx.x + blockIdx.x*blockDim.x;
|
||||
int offset = blockDim.x*gridDim.x;
|
||||
for(int i=tid; i<N; i+=offset) {
|
||||
T c = contrib[i];
|
||||
T a = acum[i];
|
||||
if(OP == ncclSum) {
|
||||
acum[i] = a+c;
|
||||
} else if(OP == ncclProd) {
|
||||
acum[i] = a*c;
|
||||
} else if(OP == ncclMax) {
|
||||
acum[i] = (a > c) ? a : c;
|
||||
} else if(OP == ncclMin) {
|
||||
acum[i] = (a < c) ? a : c;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#ifdef CUDA_HAS_HALF
|
||||
template<> __global__
|
||||
void accumKern<half, ncclSum>(half* acum, const half* contrib, int N) {
|
||||
int tid = threadIdx.x + blockIdx.x*blockDim.x;
|
||||
int offset = blockDim.x*gridDim.x;
|
||||
for(int i=tid; i<N; i+=offset) {
|
||||
float c = __half2float(contrib[i]);
|
||||
float a = __half2float(acum[i]);
|
||||
acum[i] = __float2half( a + c );
|
||||
}
|
||||
}
|
||||
|
||||
template<> __global__
|
||||
void accumKern<half, ncclProd>(half* acum, const half* contrib, int N) {
|
||||
int tid = threadIdx.x + blockIdx.x*blockDim.x;
|
||||
int offset = blockDim.x*gridDim.x;
|
||||
for(int i=tid; i<N; i+=offset) {
|
||||
float c = __half2float(contrib[i]);
|
||||
float a = __half2float(acum[i]);
|
||||
acum[i] = __float2half( a * c );
|
||||
}
|
||||
}
|
||||
|
||||
template<> __global__
|
||||
void accumKern<half, ncclMax>(half* acum, const half* contrib, int N) {
|
||||
int tid = threadIdx.x + blockIdx.x*blockDim.x;
|
||||
int offset = blockDim.x*gridDim.x;
|
||||
for(int i=tid; i<N; i+=offset) {
|
||||
float c = __half2float(contrib[i]);
|
||||
float a = __half2float(acum[i]);
|
||||
acum[i] = __float2half( (a>c) ? a : c );
|
||||
}
|
||||
}
|
||||
|
||||
template<> __global__
|
||||
void accumKern<half, ncclMin>(half* acum, const half* contrib, int N) {
|
||||
int tid = threadIdx.x + blockIdx.x*blockDim.x;
|
||||
int offset = blockDim.x*gridDim.x;
|
||||
for(int i=tid; i<N; i+=offset) {
|
||||
float c = __half2float(contrib[i]);
|
||||
float a = __half2float(acum[i]);
|
||||
acum[i] = __float2half( (a<c) ? a : c );
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
template<typename T>
|
||||
void Accumulate(T* dest, const T* contrib, int N, ncclRedOp_t op) {
|
||||
|
||||
T* devdest;
|
||||
CUDACHECK(cudaHostRegister(dest, N*sizeof(T), 0));
|
||||
CUDACHECK(cudaHostGetDevicePointer(&devdest, dest, 0));
|
||||
switch(op) {
|
||||
case ncclSum: accumKern<T, ncclSum> <<<256,256>>>(devdest, contrib, N); break;
|
||||
case ncclProd: accumKern<T, ncclProd><<<256,256>>>(devdest, contrib, N); break;
|
||||
case ncclMax: accumKern<T, ncclMax> <<<256,256>>>(devdest, contrib, N); break;
|
||||
case ncclMin: accumKern<T, ncclMin> <<<256,256>>>(devdest, contrib, N); break;
|
||||
default:
|
||||
printf("Unknown reduction operation.\n");
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
CUDACHECK(cudaHostUnregister(dest));
|
||||
}
|
||||
|
||||
template<typename T> __device__
|
||||
double absDiff(T a, T b) {
|
||||
return fabs((double)(b - a));
|
||||
}
|
||||
|
||||
#ifdef CUDA_HAS_HALF
|
||||
template<> __device__
|
||||
double absDiff<half>(half a, half b) {
|
||||
float x = __half2float(a);
|
||||
float y = __half2float(b);
|
||||
return fabs((double)(y-x));
|
||||
}
|
||||
#endif
|
||||
|
||||
template<typename T, int BSIZE> __global__
|
||||
void deltaKern(const T* A, const T* B, int N, double* max) {
|
||||
__shared__ double temp[BSIZE];
|
||||
int tid = threadIdx.x;
|
||||
double locmax = 0.0;
|
||||
for(int i=tid; i<N; i+=blockDim.x) {
|
||||
|
||||
double delta = absDiff(A[i], B[i]);
|
||||
if( delta > locmax )
|
||||
locmax = delta;
|
||||
}
|
||||
|
||||
temp[tid] = locmax;
|
||||
for(int stride = BSIZE/2; stride > 1; stride>>=1) {
|
||||
__syncthreads();
|
||||
if( tid < stride )
|
||||
temp[tid] = temp[tid] > temp[tid+stride] ? temp[tid] : temp[tid+stride];
|
||||
}
|
||||
__syncthreads();
|
||||
if( threadIdx.x == 0)
|
||||
*max = temp[0] > temp[1] ? temp[0] : temp[1];
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
double CheckDelta(const T* results, const T* expected, int N) {
|
||||
T* devexp;
|
||||
double maxerr;
|
||||
double* devmax;
|
||||
CUDACHECK(cudaHostRegister((void*)expected, N*sizeof(T), 0));
|
||||
CUDACHECK(cudaHostGetDevicePointer((void**)&devexp, (void*)expected, 0));
|
||||
CUDACHECK(cudaHostRegister((void*)&maxerr, sizeof(double), 0));
|
||||
CUDACHECK(cudaHostGetDevicePointer(&devmax, &maxerr, 0));
|
||||
deltaKern<T, 512><<<1, 512>>>(results, devexp, N, devmax);
|
||||
CUDACHECK(cudaHostUnregister(&maxerr));
|
||||
CUDACHECK(cudaHostUnregister((void*)devexp));
|
||||
return maxerr;
|
||||
}
|
||||
|
||||
|
||||
std::string TypeName(const ncclDataType_t type) {
|
||||
switch (type) {
|
||||
case ncclChar: return "char";
|
||||
case ncclInt: return "int";
|
||||
#ifdef CUDA_HAS_HALF
|
||||
case ncclHalf: return "half";
|
||||
#endif
|
||||
case ncclFloat: return "float";
|
||||
case ncclDouble: return "double";
|
||||
case ncclInt64: return "int64";
|
||||
case ncclUint64: return "uint64";
|
||||
default: return "unknown";
|
||||
}
|
||||
}
|
||||
|
||||
std::string OperationName(const ncclRedOp_t op) {
|
||||
switch (op) {
|
||||
case ncclSum: return "sum";
|
||||
case ncclProd: return "prod";
|
||||
case ncclMax: return "max";
|
||||
case ncclMin: return "min";
|
||||
default: return "unknown";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
#endif // SRC_TEST_UTILITIES_H_
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele