AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

gtest: add scatter, gather and all to all unit tests

Просмотр исходного кода
Этот коммит содержится в:
Wenkai Du
2020-04-28 16:26:42 -07:00
родитель e98891d039
Коммит fee1a20b74
17 изменённых файлов: 380 добавлений и 24 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
test/CMakeLists.txt
+3
Просмотреть файл
@@ -48,6 +48,9 @@ if(BUILD_TESTS)
test_CombinedCalls.cpp
test_AllReduceAbort.cpp
test_BroadcastAbort.cpp
test_Scatter.cpp
test_Gather.cpp
test_AllToAll.cpp
)
add_executable(UnitTests ${TEST_SOURCES})
test/CorrectnessTest.hpp
+77 -15
Просмотреть файл
@@ -16,8 +16,13 @@
#define HIP_CALL(x) ASSERT_EQ(x, hipSuccess)
#define NCCL_CALL(x) ASSERT_EQ(x, ncclSuccess)
#define MAX_ENV_TOKENS 16
namespace CorrectnessTests
{
typedef enum { ncclCollBroadcast, ncclCollReduce, ncclCollAllGather, ncclCollReduceScatter, ncclCollAllReduce, ncclCollGather, ncclCollScatter, ncclCollAllToAll, ncclCollSendRecv } ncclFunc_t;
typedef enum { ncclInputBuffer, ncclOutputBuffer } ncclBufferType_t;
// Performs the various basic reduction operations
template <typename T>
T ReduceOp(ncclRedOp_t const op, T const A, T const B)
@@ -62,6 +67,7 @@ namespace CorrectnessTests
size_t numElements; // Number of elements per array
ncclDataType_t dataType; // Data type of each input/output pointer
bool inPlace; // Whether or not output pointers are same as input pointers
ncclFunc_t function; // Buffer sizes are different in case of gather, scatter and all to all
std::vector<void *> inputs; // Input pointers (1 per device)
std::vector<void *> outputs; // Output pointers (1 per device)
@@ -73,33 +79,42 @@ namespace CorrectnessTests
return numElements * DataTypeToBytes(dataType);
}
size_t NumBytes(ncclBufferType_t bufferType) const
{
if ((function == ncclCollGather && (bufferType == ncclOutputBuffer || inPlace == true)) ||
(function == ncclCollScatter && bufferType == ncclInputBuffer) ||
function == ncclCollAllToAll)
return numElements * DataTypeToBytes(dataType) * numDevices;
return numElements * DataTypeToBytes(dataType);
}
void Initialize(int const numDevices_,
size_t const numElements_,
ncclDataType_t const dataType_,
bool const inPlace_)
bool const inPlace_,
ncclFunc_t const func_ = ncclCollBroadcast)
{
numDevices = numDevices_;
numElements = numElements_;
dataType = dataType_;
inPlace = inPlace_;
function = func_;
inputs.resize(numDevices);
outputs.resize(numDevices);
expected.resize(numDevices);
// Allocate per-device memory
size_t const numBytes = NumBytes();
for (int i = 0; i < numDevices; i++)
{
HIP_CALL(hipSetDevice(i));
HIP_CALL(hipMalloc((void **)&inputs[i], numBytes));
HIP_CALL(hipMalloc((void **)&inputs[i], NumBytes(ncclInputBuffer)));
if (inPlace)
outputs[i] = inputs[i];
else
HIP_CALL(hipMalloc((void **)&outputs[i], numBytes));
HIP_CALL(hipMalloc((void **)&outputs[i], NumBytes(ncclOutputBuffer)));
expected[i] = malloc(numBytes);
expected[i] = malloc(NumBytes(ncclOutputBuffer));
}
}
@@ -148,7 +163,8 @@ namespace CorrectnessTests
ncclDataType_t /* dataType */,
size_t /* numElements */,
int /* numDevices */,
bool /* inPlace */> TestTuple;
bool /* inPlace */,
const char* /* envVals */> TestTuple;
// Base class for each collective test
// - Each test is instantiated with a different TestTuple
@@ -167,7 +183,28 @@ namespace CorrectnessTests
}
// Make the test tuple parameters accessible
std::tie(op, dataType, numElements, numDevices, inPlace) = GetParam();
std::tie(op, dataType, numElements, numDevices, inPlace, envVals) = GetParam();
envString = 0;
numTokens = 0;
if (strcmp(envVals, "")) {
// enable RCCL env vars testing
setenv("RCCL_TEST_ENV_VARS", "ENABLE", 1);
envString = strdup(envVals);
tokens[numTokens] = strtok(envString, "=, ");
numTokens++;
while (tokens[numTokens-1] != NULL && numTokens < MAX_ENV_TOKENS)
tokens[numTokens++] = strtok(NULL, "=, ");
for (int i = 0; i < numTokens/2; i++) {
char *val = getenv(tokens[i*2]);
if (val)
savedEnv[i] = strdup(val);
else
savedEnv[i] = 0;
setenv(tokens[i*2], tokens[i*2+1], 1);
fprintf(stdout, "[ ] setting environmental variable %s to %s\n", tokens[i*2], getenv(tokens[i*2]));
}
}
// Collect the number of available GPUs
HIP_CALL(hipGetDeviceCount(&numDevicesAvailable));
@@ -207,11 +244,26 @@ namespace CorrectnessTests
NCCL_CALL(ncclCommDestroy(comms[i]));
HIP_CALL(hipStreamDestroy(streams[i]));
}
// Restore env vars after tests
for (int i = 0; i < numTokens/2; i++) {
if (savedEnv[i]) {
setenv(tokens[i*2], savedEnv[i], 1);
fprintf(stdout, "[ ] restored environmental variable %s to %s\n", tokens[i*2], getenv(tokens[i*2]));
free(savedEnv[i]);
}
else {
unsetenv(tokens[i*2]);
fprintf(stdout, "[ ] removed environmental variable %s\n", tokens[i*2]);
}
}
// Cleanup
unsetenv("RCCL_TEST_ENV_VARS");
free(envString);
}
void FillDatasetWithPattern(Dataset& dataset)
{
int8_t* arrayI1 = (int8_t *)malloc(dataset.NumBytes());
int8_t* arrayI1 = (int8_t *)malloc(dataset.NumBytes(ncclInputBuffer));
uint8_t* arrayU1 = (uint8_t *)arrayI1;
int32_t* arrayI4 = (int32_t *)arrayI1;
uint32_t* arrayU4 = (uint32_t *)arrayI1;
@@ -229,7 +281,7 @@ namespace CorrectnessTests
// - Sticking with floating points values that are perfectly representable
for (int i = 0; i < dataset.numDevices; i++)
{
for (int j = 0; j < dataset.numElements; j++)
for (int j = 0; j < dataset.NumBytes(ncclInputBuffer)/DataTypeToBytes(dataset.dataType); j++)
{
int valueI = (i + j) % 6;
float valueF = (float)valueI;
@@ -252,11 +304,11 @@ namespace CorrectnessTests
}
HIP_CALL(hipSetDevice(i));
HIP_CALL(hipMemcpy(dataset.inputs[i], arrayI1, dataset.NumBytes(), hipMemcpyHostToDevice));
HIP_CALL(hipMemcpy(dataset.inputs[i], arrayI1, dataset.NumBytes(ncclInputBuffer), hipMemcpyHostToDevice));
// Fills output data[i][j] with 0 (if not inplace)
if (!dataset.inPlace)
HIP_CALL(hipMemset(dataset.outputs[i], 0, dataset.NumBytes()));
HIP_CALL(hipMemset(dataset.outputs[i], 0, dataset.NumBytes(ncclOutputBuffer)));
}
free(arrayI1);
@@ -272,9 +324,9 @@ namespace CorrectnessTests
}
}
void ValidateResults(Dataset const& dataset) const
void ValidateResults(Dataset const& dataset, int root = 0) const
{
int8_t* outputI1 = (int8_t *)malloc(dataset.NumBytes());
int8_t* outputI1 = (int8_t *)malloc(dataset.NumBytes(ncclOutputBuffer));
uint8_t* outputU1 = (uint8_t *)outputI1;
int32_t* outputI4 = (int32_t *)outputI1;
uint32_t* outputU4 = (uint32_t *)outputI1;
@@ -290,7 +342,10 @@ namespace CorrectnessTests
// (Each collective operation computes its own expected results)
for (int i = 0; i < dataset.numDevices && isMatch; i++)
{
HIP_CALL(hipMemcpy(outputI1, dataset.outputs[i], dataset.NumBytes(), hipMemcpyDeviceToHost));
// only output on root rank is valid for gather collective
if (dataset.function == ncclCollGather && i != root)
continue;
HIP_CALL(hipMemcpy(outputI1, dataset.outputs[i], dataset.NumBytes(ncclOutputBuffer), hipMemcpyDeviceToHost));
int8_t* expectedI1 = (int8_t *)dataset.expected[i];
uint8_t* expectedU1 = (uint8_t *)expectedI1;
@@ -358,10 +413,17 @@ namespace CorrectnessTests
size_t numElements;
int numDevices;
bool inPlace;
const char* envVals;
int numDevicesAvailable;
std::vector<ncclComm_t> comms;
std::vector<hipStream_t> streams;
// internal only
char* envString;
char* tokens[MAX_ENV_TOKENS];
int numTokens;
char* savedEnv[MAX_ENV_TOKENS/2];
};
}
test/test_AllGather.cpp
+2 -1
Просмотреть файл
@@ -109,5 +109,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false, true)));
testing::Values(false, true),
testing::Values("")));
} // namespace
test/test_AllReduce.cpp
+2 -1
Просмотреть файл
@@ -57,5 +57,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false, true)));
testing::Values(false, true),
testing::Values("")));
} // namespace
test/test_AllReduceAbort.cpp
+2 -1
Просмотреть файл
@@ -138,5 +138,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2, 4),
// In-place or not
testing::Values(false)));
testing::Values(false),
testing::Values("")));
} // namespace
test/test_AllToAll.cpp
+66
Просмотреть файл
@@ -0,0 +1,66 @@
/*************************************************************************
* Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#include "test_AllToAll.hpp"
namespace CorrectnessTests
{
TEST_P(AllToAllCorrectnessTest, Correctness)
{
if (numDevices > numDevicesAvailable) return;
// Allocate data
Dataset dataset;
dataset.Initialize(numDevices, numElements, dataType, inPlace, ncclCollAllToAll);
// Prepare input / output / expected results
FillDatasetWithPattern(dataset);
ComputeExpectedResults(dataset);
// Launch the reduction (1 thread per GPU)
ncclGroupStart();
for (int i = 0; i < numDevices; i++)
{
ncclAllToAll(dataset.inputs[i],
dataset.outputs[i],
numElements, dataType,
comms[i], streams[i]);
}
ncclGroupEnd();
// Wait for reduction to complete
Synchronize();
// Check results
ValidateResults(dataset);
dataset.Release();
}
INSTANTIATE_TEST_CASE_P(AllToAllCorrectnessSweep,
AllToAllCorrectnessTest,
testing::Combine(
// Reduction operator is not used
testing::Values(ncclSum),
// Data types
testing::Values(ncclInt8,
ncclUint8,
ncclInt32,
ncclUint32,
ncclInt64,
ncclUint64,
//ncclFloat16,
ncclFloat32,
ncclFloat64,
ncclBfloat16),
// Number of elements
testing::Values(1024, 1048576),
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false),
testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=0", "RCCL_ALLTOALL_KERNEL_DISABLE=1")));
} // namespace
test/test_AllToAll.hpp
+26
Просмотреть файл
@@ -0,0 +1,26 @@
/*************************************************************************
* Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#ifndef TEST_ALLTOALL_HPP
#define TEST_ALLTOALL_HPP
#include "CorrectnessTest.hpp"
namespace CorrectnessTests
{
class AllToAllCorrectnessTest : public CorrectnessTest
{
public:
static void ComputeExpectedResults(Dataset& dataset)
{
for (int i = 0; i < dataset.numDevices; i++)
for (int j = 0; j < dataset.numDevices; j++)
HIP_CALL(hipMemcpy((int8_t *)dataset.expected[i]+dataset.NumBytes()*j, (int8_t *)dataset.inputs[j]+dataset.NumBytes()*i,
dataset.NumBytes(), hipMemcpyDeviceToHost));
}
};
}
#endif
test/test_Broadcast.cpp
+2 -1
Просмотреть файл
@@ -65,5 +65,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false, true)));
testing::Values(false, true),
testing::Values("")));
} // namespace
test/test_BroadcastAbort.cpp
+2 -1
Просмотреть файл
@@ -141,5 +141,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2, 4),
// In-place or not
testing::Values(false)));
testing::Values(false),
testing::Values("")));
} // namespace
test/test_CombinedCalls.cpp
+2 -1
Просмотреть файл
@@ -95,5 +95,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false, true)));
testing::Values(false, true),
testing::Values("")));
} // namespace
test/test_Gather.cpp
+70
Просмотреть файл
@@ -0,0 +1,70 @@
/*************************************************************************
* Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#include "test_Gather.hpp"
namespace CorrectnessTests
{
TEST_P(GatherCorrectnessTest, Correctness)
{
if (numDevices > numDevicesAvailable) return;
// Allocate data
Dataset dataset;
dataset.Initialize(numDevices, numElements, dataType, inPlace, ncclCollGather);
// Test each possible root
for (int root = 0; root < numDevices; root++)
{
// Prepare input / output / expected results
FillDatasetWithPattern(dataset);
ComputeExpectedResults(dataset, root);
// Launch the reduction (1 thread per GPU)
ncclGroupStart();
for (int i = 0; i < numDevices; i++)
{
ncclGather(dataset.inputs[i],
dataset.outputs[i],
numElements, dataType,
root, comms[i], streams[i]);
}
ncclGroupEnd();
// Wait for reduction to complete
Synchronize();
// Check results
ValidateResults(dataset, root);
}
dataset.Release();
}
INSTANTIATE_TEST_CASE_P(GatherCorrectnessSweep,
GatherCorrectnessTest,
testing::Combine(
// Reduction operator is not used
testing::Values(ncclSum),
// Data types
testing::Values(ncclInt8,
ncclUint8,
ncclInt32,
ncclUint32,
ncclInt64,
ncclUint64,
//ncclFloat16,
ncclFloat32,
ncclFloat64,
ncclBfloat16),
// Number of elements
testing::Values(1024, 1048576),
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false),
testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=0", "RCCL_ALLTOALL_KERNEL_DISABLE=1")));
} // namespace
test/test_Gather.hpp
+25
Просмотреть файл
@@ -0,0 +1,25 @@
/*************************************************************************
* Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#ifndef TEST_GATHER_HPP
#define TEST_GATHER_HPP
#include "CorrectnessTest.hpp"
namespace CorrectnessTests
{
class GatherCorrectnessTest : public CorrectnessTest
{
public:
static void ComputeExpectedResults(Dataset& dataset, int const root)
{
for (int i = 0; i < dataset.numDevices; i++)
HIP_CALL(hipMemcpy((int8_t *)dataset.expected[root]+dataset.NumBytes()*i, dataset.inputs[i],
dataset.NumBytes(), hipMemcpyDeviceToHost));
}
};
}
#endif
test/test_GroupCalls.cpp
+2 -1
Просмотреть файл
@@ -115,5 +115,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false, true)));
testing::Values(false, true),
testing::Values("")));
} // namespace
test/test_Reduce.cpp
+2 -1
Просмотреть файл
@@ -65,5 +65,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false, true)));
testing::Values(false, true),
testing::Values("")));
} // namespace
test/test_ReduceScatter.cpp
+2 -1
Просмотреть файл
@@ -63,5 +63,6 @@ namespace CorrectnessTests
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false, true)));
testing::Values(false, true),
testing::Values("")));
} // namespace
test/test_Scatter.cpp
+70
Просмотреть файл
@@ -0,0 +1,70 @@
/*************************************************************************
* Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#include "test_Scatter.hpp"
namespace CorrectnessTests
{
TEST_P(ScatterCorrectnessTest, Correctness)
{
if (numDevices > numDevicesAvailable) return;
// Allocate data
Dataset dataset;
dataset.Initialize(numDevices, numElements, dataType, inPlace, ncclCollScatter);
// Test each possible root
for (int root = 0; root < numDevices; root++)
{
// Prepare input / output / expected results
FillDatasetWithPattern(dataset);
ComputeExpectedResults(dataset, root);
// Launch the reduction (1 thread per GPU)
ncclGroupStart();
for (int i = 0; i < numDevices; i++)
{
ncclScatter(dataset.inputs[i],
dataset.outputs[i],
numElements, dataType,
root, comms[i], streams[i]);
}
ncclGroupEnd();
// Wait for reduction to complete
Synchronize();
// Check results
ValidateResults(dataset);
}
dataset.Release();
}
INSTANTIATE_TEST_CASE_P(ScatterCorrectnessSweep,
ScatterCorrectnessTest,
testing::Combine(
// Reduction operator is not used
testing::Values(ncclSum),
// Data types
testing::Values(ncclInt8,
ncclUint8,
ncclInt32,
ncclUint32,
ncclInt64,
ncclUint64,
//ncclFloat16,
ncclFloat32,
ncclFloat64,
ncclBfloat16),
// Number of elements
testing::Values(1024, 1048576),
// Number of devices
testing::Values(2,3,4),
// In-place or not
testing::Values(false),
testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=0", "RCCL_ALLTOALL_KERNEL_DISABLE=1")));
} // namespace
test/test_Scatter.hpp
+25
Просмотреть файл
@@ -0,0 +1,25 @@
/*************************************************************************
* Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#ifndef TEST_SCATTER_HPP
#define TEST_SCATTER_HPP
#include "CorrectnessTest.hpp"
namespace CorrectnessTests
{
class ScatterCorrectnessTest : public CorrectnessTest
{
public:
static void ComputeExpectedResults(Dataset& dataset, int const root)
{
for (int i = 0; i < dataset.numDevices; i++)
HIP_CALL(hipMemcpy(dataset.expected[i], (int8_t *)dataset.inputs[root]+dataset.NumBytes()*i,
dataset.NumBytes(), hipMemcpyDeviceToHost));
}
};
}
#endif