From fee1a20b74f52c9c20c1b0a405aaa9672dbc83a2 Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Tue, 28 Apr 2020 16:26:42 -0700 Subject: [PATCH] gtest: add scatter, gather and all to all unit tests --- test/CMakeLists.txt | 3 ++ test/CorrectnessTest.hpp | 92 ++++++++++++++++++++++++++++++------ test/test_AllGather.cpp | 3 +- test/test_AllReduce.cpp | 3 +- test/test_AllReduceAbort.cpp | 3 +- test/test_AllToAll.cpp | 66 ++++++++++++++++++++++++++ test/test_AllToAll.hpp | 26 ++++++++++ test/test_Broadcast.cpp | 3 +- test/test_BroadcastAbort.cpp | 3 +- test/test_CombinedCalls.cpp | 3 +- test/test_Gather.cpp | 70 +++++++++++++++++++++++++++ test/test_Gather.hpp | 25 ++++++++++ test/test_GroupCalls.cpp | 3 +- test/test_Reduce.cpp | 3 +- test/test_ReduceScatter.cpp | 3 +- test/test_Scatter.cpp | 70 +++++++++++++++++++++++++++ test/test_Scatter.hpp | 25 ++++++++++ 17 files changed, 380 insertions(+), 24 deletions(-) create mode 100644 test/test_AllToAll.cpp create mode 100644 test/test_AllToAll.hpp create mode 100644 test/test_Gather.cpp create mode 100644 test/test_Gather.hpp create mode 100644 test/test_Scatter.cpp create mode 100644 test/test_Scatter.hpp diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 547553b2a0..825e60d4dc 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -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}) diff --git a/test/CorrectnessTest.hpp b/test/CorrectnessTest.hpp index ba6cb6090c..94758c68d7 100644 --- a/test/CorrectnessTest.hpp +++ b/test/CorrectnessTest.hpp @@ -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 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 inputs; // Input pointers (1 per device) std::vector 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 comms; std::vector streams; + + // internal only + char* envString; + char* tokens[MAX_ENV_TOKENS]; + int numTokens; + char* savedEnv[MAX_ENV_TOKENS/2]; }; } diff --git a/test/test_AllGather.cpp b/test/test_AllGather.cpp index c0e92a8ee5..558e927ffb 100644 --- a/test/test_AllGather.cpp +++ b/test/test_AllGather.cpp @@ -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 diff --git a/test/test_AllReduce.cpp b/test/test_AllReduce.cpp index 8d17339961..41f1af5a8b 100644 --- a/test/test_AllReduce.cpp +++ b/test/test_AllReduce.cpp @@ -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 diff --git a/test/test_AllReduceAbort.cpp b/test/test_AllReduceAbort.cpp index 69edb71882..d84f36f8ad 100644 --- a/test/test_AllReduceAbort.cpp +++ b/test/test_AllReduceAbort.cpp @@ -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 diff --git a/test/test_AllToAll.cpp b/test/test_AllToAll.cpp new file mode 100644 index 0000000000..2c1d29da1b --- /dev/null +++ b/test/test_AllToAll.cpp @@ -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 diff --git a/test/test_AllToAll.hpp b/test/test_AllToAll.hpp new file mode 100644 index 0000000000..104daba05d --- /dev/null +++ b/test/test_AllToAll.hpp @@ -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 diff --git a/test/test_Broadcast.cpp b/test/test_Broadcast.cpp index 3639ef1c71..85b087a163 100644 --- a/test/test_Broadcast.cpp +++ b/test/test_Broadcast.cpp @@ -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 diff --git a/test/test_BroadcastAbort.cpp b/test/test_BroadcastAbort.cpp index feaa09b004..e14c8f159c 100644 --- a/test/test_BroadcastAbort.cpp +++ b/test/test_BroadcastAbort.cpp @@ -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 diff --git a/test/test_CombinedCalls.cpp b/test/test_CombinedCalls.cpp index 6c17ea6f57..35d61e82f3 100644 --- a/test/test_CombinedCalls.cpp +++ b/test/test_CombinedCalls.cpp @@ -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 diff --git a/test/test_Gather.cpp b/test/test_Gather.cpp new file mode 100644 index 0000000000..4324434e9e --- /dev/null +++ b/test/test_Gather.cpp @@ -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 diff --git a/test/test_Gather.hpp b/test/test_Gather.hpp new file mode 100644 index 0000000000..f75cfa9b2a --- /dev/null +++ b/test/test_Gather.hpp @@ -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 diff --git a/test/test_GroupCalls.cpp b/test/test_GroupCalls.cpp index 365f72067f..cc2bd7dcd2 100644 --- a/test/test_GroupCalls.cpp +++ b/test/test_GroupCalls.cpp @@ -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 diff --git a/test/test_Reduce.cpp b/test/test_Reduce.cpp index 484e1b02e3..c181ec90ac 100644 --- a/test/test_Reduce.cpp +++ b/test/test_Reduce.cpp @@ -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 diff --git a/test/test_ReduceScatter.cpp b/test/test_ReduceScatter.cpp index 51e510a403..014ecd0ba4 100644 --- a/test/test_ReduceScatter.cpp +++ b/test/test_ReduceScatter.cpp @@ -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 diff --git a/test/test_Scatter.cpp b/test/test_Scatter.cpp new file mode 100644 index 0000000000..dca86b1ff3 --- /dev/null +++ b/test/test_Scatter.cpp @@ -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 diff --git a/test/test_Scatter.hpp b/test/test_Scatter.hpp new file mode 100644 index 0000000000..a7e695c28b --- /dev/null +++ b/test/test_Scatter.hpp @@ -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