From 57ac9a8a9339acebe4a89ff44e6702df47e94a23 Mon Sep 17 00:00:00 2001 From: Gilbert Lee Date: Fri, 17 May 2019 23:57:11 +0000 Subject: [PATCH] Adding support for alignment tests via sub-datasets Added sample alignment test for AllGather Datasets no longer free memory on destruction so Release() must be used [ROCm/rccl commit: a50c85285135e9b0ba610f1d68623ff2993ce370] --- projects/rccl/test/CorrectnessTest.hpp | 33 ++++++++++++++- projects/rccl/test/test_AllGather.cpp | 49 +++++++++++++++++++++++ projects/rccl/test/test_AllReduce.cpp | 2 + projects/rccl/test/test_Broadcast.cpp | 2 + projects/rccl/test/test_GroupCalls.cpp | 1 + projects/rccl/test/test_Reduce.cpp | 2 + projects/rccl/test/test_ReduceScatter.cpp | 2 + 7 files changed, 90 insertions(+), 1 deletion(-) diff --git a/projects/rccl/test/CorrectnessTest.hpp b/projects/rccl/test/CorrectnessTest.hpp index a4dbd567ac..d0a4680ef3 100644 --- a/projects/rccl/test/CorrectnessTest.hpp +++ b/projects/rccl/test/CorrectnessTest.hpp @@ -60,6 +60,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 + std::vector inputs; // Input pointers (1 per device) std::vector outputs; // Output pointers (1 per device) // May be identical to input pointers for in-place tests @@ -100,7 +101,8 @@ namespace CorrectnessTests } } - ~Dataset() + // Explicit memory release to avoid double-free from subDatasets + void Release() { for (int i = 0; i < outputs.size(); i++) { @@ -108,6 +110,35 @@ namespace CorrectnessTests hipFree(inputs[i]); free(expected[i]); } + + outputs.clear(); + } + + // Creates a dataset by pointing to an existing dataset + // Primarily to allow for testing with different starting byte-alignments + void ExtractSubDataset(size_t const startElement, + size_t const lastElement, + Dataset& subDataset) + { + ASSERT_LE(startElement, lastElement); + ASSERT_LT(lastElement, numElements); + + subDataset.numDevices = numDevices; + subDataset.numElements = lastElement - startElement + 1; + subDataset.dataType = dataType; + subDataset.inPlace = inPlace; + + subDataset.inputs.resize(numDevices); + subDataset.outputs.resize(numDevices); + subDataset.expected.resize(numDevices); + + size_t const byteOffset = (startElement * DataTypeToBytes(dataType)); + for (int i = 0; i < numDevices; i++) + { + subDataset.inputs[i] = (int8_t *)inputs[i] + byteOffset; + subDataset.outputs[i] = (int8_t *)outputs[i] + byteOffset; + subDataset.expected[i] = (int8_t *)expected[i] + byteOffset; + } } }; diff --git a/projects/rccl/test/test_AllGather.cpp b/projects/rccl/test/test_AllGather.cpp index c6e79fe9c8..2727514186 100644 --- a/projects/rccl/test/test_AllGather.cpp +++ b/projects/rccl/test/test_AllGather.cpp @@ -36,8 +36,57 @@ namespace CorrectnessTests // Check results ValidateResults(dataset); + dataset.Release(); } + TEST_P(AllGatherCorrectnessTest, Alignment) + { + if (numDevices > numDevicesAvailable) return; + if (numElements % numDevices != 0) return; + + // Allocate dataset + Dataset dataset; + dataset.Initialize(numDevices, numElements, dataType, inPlace); + + // Loop over several offsets (so that device pointers are not aligned) + for (int firstElement = 1; firstElement <= 11; firstElement += 2) + { + if (firstElement < numElements) + { + // Select last element so that total number of elements is multiple of numDevices + int const lastElement = firstElement + ((numElements - firstElement) / numDevices) * numDevices - 1; + if (lastElement >= numElements) break; + + Dataset subDataset; + dataset.ExtractSubDataset(firstElement, lastElement, subDataset); + + // Compute reference results for sub-dataset + FillDatasetWithPattern(subDataset); + ComputeExpectedResults(subDataset); + + size_t const byteCount = subDataset.NumBytes() / subDataset.numDevices; + size_t const sendCount = subDataset.numElements / subDataset.numDevices; + + // Launch the reduction (1 thread per GPU) + #pragma omp parallel for num_threads(numDevices) + for (int i = 0; i < numDevices; i++) + { + ncclAllGather((int8_t *)subDataset.inputs[i] + (i * byteCount), + subDataset.outputs[i], sendCount, + dataType, comms[i], streams[i]); + } + + // Wait for reduction to complete + Synchronize(); + + // Check results + ValidateResults(subDataset); + } + } + dataset.Release(); + } + + INSTANTIATE_TEST_CASE_P(AllGatherCorrectnessSweep, AllGatherCorrectnessTest, testing::Combine( diff --git a/projects/rccl/test/test_AllReduce.cpp b/projects/rccl/test/test_AllReduce.cpp index f77651c84e..0fb7474d0e 100644 --- a/projects/rccl/test/test_AllReduce.cpp +++ b/projects/rccl/test/test_AllReduce.cpp @@ -32,6 +32,8 @@ namespace CorrectnessTests // Check results ValidateResults(dataset); + + dataset.Release(); } INSTANTIATE_TEST_CASE_P(AllReduceCorrectnessSweep, diff --git a/projects/rccl/test/test_Broadcast.cpp b/projects/rccl/test/test_Broadcast.cpp index 2f2a091a6d..3ed6964785 100644 --- a/projects/rccl/test/test_Broadcast.cpp +++ b/projects/rccl/test/test_Broadcast.cpp @@ -41,6 +41,8 @@ namespace CorrectnessTests // Check results ValidateResults(dataset); } + + dataset.Release(); } INSTANTIATE_TEST_CASE_P(BroadcastCorrectnessSweep, diff --git a/projects/rccl/test/test_GroupCalls.cpp b/projects/rccl/test/test_GroupCalls.cpp index d713588c13..de1ad0bd76 100644 --- a/projects/rccl/test/test_GroupCalls.cpp +++ b/projects/rccl/test/test_GroupCalls.cpp @@ -92,6 +92,7 @@ namespace CorrectnessTests for (int i = 0; i < 5; i++) { ValidateResults(datasets[i]); + datasets[i].Release(); } } diff --git a/projects/rccl/test/test_Reduce.cpp b/projects/rccl/test/test_Reduce.cpp index bf6c2164e5..dfca79ccf4 100644 --- a/projects/rccl/test/test_Reduce.cpp +++ b/projects/rccl/test/test_Reduce.cpp @@ -40,6 +40,8 @@ namespace CorrectnessTests // Check results ValidateResults(dataset); } + + dataset.Release(); } INSTANTIATE_TEST_CASE_P(ReduceCorrectnessSweep, diff --git a/projects/rccl/test/test_ReduceScatter.cpp b/projects/rccl/test/test_ReduceScatter.cpp index 567ce1fef1..d55b514689 100644 --- a/projects/rccl/test/test_ReduceScatter.cpp +++ b/projects/rccl/test/test_ReduceScatter.cpp @@ -39,6 +39,8 @@ namespace CorrectnessTests // Check results ValidateResults(dataset); + + dataset.Release(); } INSTANTIATE_TEST_CASE_P(ReduceScatterCorrectnessSweep,