diff --git a/projects/rccl/.jenkins/common.groovy b/projects/rccl/.jenkins/common.groovy index fe5a245be9..24cefcdb47 100644 --- a/projects/rccl/.jenkins/common.groovy +++ b/projects/rccl/.jenkins/common.groovy @@ -14,7 +14,7 @@ def runCompileCommand(platform, project, jobName) platform.runCommand(this,command) } -def runTestCommand (platform, project, gfilter) +def runTestCommand (platform, project, gfilter, envars) { String sudo = auxiliary.sudo(platform.jenkinsLabel) @@ -23,8 +23,8 @@ def runTestCommand (platform, project, gfilter) cd ${project.paths.project_build_prefix}/build/release/test ${sudo} ulimit -l unlimited ulimit -a - ${sudo} UT_PROCESS_MASK=1 RCCL_ENABLE_SIGNALHANDLER=0 NCCL_DEBUG=INFO HSA_FORCE_FINE_GRAIN_PCIE=1 ./rccl-UnitTests --gtest_filter=${gfilter} --gtest_output=xml --gtest_color=yes - ${sudo} UT_PROCESS_MASK=2 RCCL_ENABLE_SIGNALHANDLER=0 NCCL_DEBUG=INFO HSA_FORCE_FINE_GRAIN_PCIE=1 ./rccl-UnitTests --gtest_filter=${gfilter} --gtest_output=xml --gtest_color=yes + ${sudo} ${envars} UT_PROCESS_MASK=1 RCCL_ENABLE_SIGNALHANDLER=0 NCCL_DEBUG=INFO HSA_FORCE_FINE_GRAIN_PCIE=1 ./rccl-UnitTests --gtest_filter=${gfilter} --gtest_output=xml --gtest_color=yes + ${sudo} ${envars} UT_PROCESS_MASK=2 RCCL_ENABLE_SIGNALHANDLER=0 NCCL_DEBUG=INFO HSA_FORCE_FINE_GRAIN_PCIE=1 ./rccl-UnitTests --gtest_filter=${gfilter} --gtest_output=xml --gtest_color=yes """ platform.runCommand(this, command) diff --git a/projects/rccl/.jenkins/extended.groovy b/projects/rccl/.jenkins/extended.groovy index ccc326dae2..7e2f2f123e 100644 --- a/projects/rccl/.jenkins/extended.groovy +++ b/projects/rccl/.jenkins/extended.groovy @@ -37,7 +37,7 @@ def runCI = { platform, project-> - commonGroovy.runTestCommand(platform, project, "*") + commonGroovy.runTestCommand(platform, project, "-*Graph", "") } def packageCommand = diff --git a/projects/rccl/.jenkins/precheckin.groovy b/projects/rccl/.jenkins/precheckin.groovy index c1d2427c9f..b093ba41d8 100644 --- a/projects/rccl/.jenkins/precheckin.groovy +++ b/projects/rccl/.jenkins/precheckin.groovy @@ -37,7 +37,7 @@ def runCI = { platform, project-> - commonGroovy.runTestCommand(platform, project, "-*ManagedMem") + commonGroovy.runTestCommand(platform, project, "-*Graph", "UT_POW2_GPUS=1") } def packageCommand = diff --git a/projects/rccl/test/AllGatherTests.cpp b/projects/rccl/test/AllGatherTests.cpp new file mode 100644 index 0000000000..edb0f31052 --- /dev/null +++ b/projects/rccl/test/AllGatherTests.cpp @@ -0,0 +1,123 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(AllGather, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllGather}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1048576, 500}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllGather, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllGather}; + std::vector const dataTypes = {ncclBfloat16, ncclFloat64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {586}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllGather, InPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllGather}; + std::vector const dataTypes = {ncclInt32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {104857, 264}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllGather, InPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllGather}; + std::vector const dataTypes = {ncclInt8, ncclInt64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {958}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllGather, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllGather}; + std::vector const dataTypes = {ncclUint8}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1039203, 2500}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllGather, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllGather}; + std::vector const dataTypes = {ncclUint32, ncclUint64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {896}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/AllGather_InPlace.cpp b/projects/rccl/test/AllGather_InPlace.cpp deleted file mode 100644 index c504ed4040..0000000000 --- a/projects/rccl/test/AllGather_InPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllGather, InPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllGather}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllGather_ManagedMem.cpp b/projects/rccl/test/AllGather_ManagedMem.cpp deleted file mode 100644 index eb265211b3..0000000000 --- a/projects/rccl/test/AllGather_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllGather, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllGather}; - std::vector const dataTypes = {ncclUint8, ncclUint32, ncclUint64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllGather_OutOfPlace.cpp b/projects/rccl/test/AllGather_OutOfPlace.cpp deleted file mode 100644 index f446c62d61..0000000000 --- a/projects/rccl/test/AllGather_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllGather, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllGather}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllReduceTests.cpp b/projects/rccl/test/AllReduceTests.cpp new file mode 100644 index 0000000000..fa771cd89f --- /dev/null +++ b/projects/rccl/test/AllReduceTests.cpp @@ -0,0 +1,217 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(AllReduce, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {393216, 384}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllReduce, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce}; + std::vector const dataTypes = {ncclFloat16, ncclFloat64}; + std::vector const redOps = {ncclMin}; + std::vector const roots = {0}; + std::vector const numElements = {12888}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllReduce, InPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce}; + std::vector const dataTypes = {ncclInt32, ncclInt8}; + std::vector const redOps = {ncclProd}; + std::vector const roots = {0}; + std::vector const numElements = {384}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllReduce, InPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce}; + std::vector const dataTypes = {ncclInt32}; + std::vector const redOps = {ncclMax}; + std::vector const roots = {0, 1}; + std::vector const numElements = {393216}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllReduce, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce}; + std::vector const dataTypes = {ncclUint8, ncclUint64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {2500}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllReduce, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce}; + std::vector const dataTypes = {ncclFloat64, ncclBfloat16}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {4314}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllReduce, DISABLED_Clique) + { + // Set clique env var prior to TestBed + setenv("RCCL_ENABLE_CLIQUE", "1", 1); + + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce}; + std::vector const dataTypes = testBed.GetAllSupportedDataTypes(); + std::vector const redOps = testBed.GetAllSupportedRedOps(); + std::vector const roots = {0}; + std::vector const numElements = {1048576, 1024}; + std::vector const inPlaceList = {false, true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false, true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + + unsetenv("RCCL_ENABLE_CLIQUE"); + } + + // This tests using custom pre-mult scalars reductions + TEST(AllReduce, PreMultScalar) + { + TestBed testBed; + + // Configuration + ncclFunc_t const funcType = ncclCollAllReduce; + std::vector const& dataTypes = {ncclFloat32}; + ncclRedOp_t const redOp = ncclSum; + std::vector const numElements = {384 * 1024, 384 * 32, 384}; + bool const inPlace = false; + bool const useManagedMem = false; + + OptionalColArgs options; + + // Terminate the test as soon as first failure occurs + bool isCorrect = true; + for (int totalRanks : testBed.ev.GetNumGpusList()) + for (int isMultiProcess : testBed.ev.GetIsMultiProcessList()) + { + int const numProcesses = isMultiProcess ? totalRanks : 1; + testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks)); + + for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) + { + ncclDataType_t const dataType = dataTypes[dataIdx]; + + // Set scalars per rank + PtrUnion scalarsPerRank; + scalarsPerRank.AllocateCpuMem(totalRanks * DataTypeToBytes(dataType)); + for (int i = 0; i < totalRanks; i++) + { + double F = i; + scalarsPerRank.Set(dataType, i, i, F); + } + int const numBytes = totalRanks * DataTypeToBytes(dataType); + memcpy(options.scalarTransport.ptr, scalarsPerRank.ptr, numBytes); + + // Test various scalar residence modes + for (int scalarMode = 0; scalarMode <= 1 && isCorrect; ++scalarMode) + { + if (testBed.ev.showNames) + INFO("%s %d-ranks AllReduce (custom-scalar Mode %d %s)\n", + isMultiProcess ? "MP" : "SP", + totalRanks, scalarMode, ncclDataTypeNames[dataType]); + + for (int i = 0; i < numElements.size() && isCorrect; ++i) + { + options.scalarMode = scalarMode; + options.redOp = redOp; + testBed.SetCollectiveArgs(funcType, dataType, + numElements[i], numElements[i], + options); + // For performance, only allocate and prepare data on largest size + if (i == 0) + { + testBed.AllocateMem(inPlace, useManagedMem); + testBed.PrepareData(); + } + testBed.ExecuteCollectives(); + testBed.ValidateResults(isCorrect); + } + testBed.DeallocateMem(); + } + } + testBed.DestroyComms(); + } + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/AllReduce_Clique.cpp b/projects/rccl/test/AllReduce_Clique.cpp deleted file mode 100644 index e1267dd524..0000000000 --- a/projects/rccl/test/AllReduce_Clique.cpp +++ /dev/null @@ -1,33 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" -#include -namespace RcclUnitTesting -{ - TEST(AllReduce, DISABLED_Clique) - { - // Set clique env var prior to TestBed - setenv("RCCL_ENABLE_CLIQUE", "1", 1); - - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllReduce}; - std::vector const dataTypes = testBed.GetAllSupportedDataTypes(); - std::vector const redOps = testBed.GetAllSupportedRedOps(); - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false, true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - - unsetenv("RCCL_ENABLE_CLIQUE"); - } -} diff --git a/projects/rccl/test/AllReduce_GroupCall.cpp b/projects/rccl/test/AllReduce_GroupCall.cpp deleted file mode 100644 index 5d10a56c83..0000000000 --- a/projects/rccl/test/AllReduce_GroupCall.cpp +++ /dev/null @@ -1,67 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllReduce, GroupCall) - { - TestBed testBed; - - // Configuration - ncclFunc_t const funcType = ncclCollAllReduce; - std::vector const& dataTypes = {ncclFloat}; - std::vector const& redOps = {ncclSum}; - std::vector const numElements = {1048576, 53327, 1024}; - bool const inPlace = false; - bool const useManagedMem = false; - int const numCollPerGroup = numElements.size(); - - OptionalColArgs options; - // This tests runs 3 collectives in the same group call - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1 && isCorrect; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - // Test either single process all GPUs, or 1 process per GPU - int const numProcesses = isMultiProcess ? totalRanks : 1; - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), numCollPerGroup); - - for (int redOpIdx = 0; redOpIdx < redOps.size() && isCorrect; ++redOpIdx) - { - options.redOp = redOps[redOpIdx]; - for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) - { - if (testBed.ev.showNames) - INFO("%s %d-ranks AllReduce %d Grouped Calls (%s-%s)\n", - isMultiProcess ? "MP" : "SP", - totalRanks, numCollPerGroup, - ncclRedOpNames[redOps[redOpIdx]], ncclDataTypeNames[dataTypes[dataIdx]]); - - // Run all element sizes in parallel as single group - for (int collIdx = 0; collIdx < numCollPerGroup; ++collIdx) - { - testBed.SetCollectiveArgs(funcType, - dataTypes[dataIdx], - numElements[collIdx], - numElements[collIdx], - options, - collIdx); - } - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - testBed.DeallocateMem(); - } - } - testBed.DestroyComms(); - } - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllReduce_InPlace.cpp b/projects/rccl/test/AllReduce_InPlace.cpp deleted file mode 100644 index e413acf870..0000000000 --- a/projects/rccl/test/AllReduce_InPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllReduce, InPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllReduce}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclFloat32}; - std::vector const redOps = {ncclSum, ncclProd}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllReduce_ManagedMem.cpp b/projects/rccl/test/AllReduce_ManagedMem.cpp deleted file mode 100644 index 1e2b1946a3..0000000000 --- a/projects/rccl/test/AllReduce_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllReduce, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllReduce}; - std::vector const dataTypes = {ncclFloat32, ncclUint8, ncclUint64}; - std::vector const redOps = {ncclSum, ncclMax}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllReduce_Msccl.cpp b/projects/rccl/test/AllReduce_Msccl.cpp deleted file mode 100644 index e4bd99cb6f..0000000000 --- a/projects/rccl/test/AllReduce_Msccl.cpp +++ /dev/null @@ -1,213 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License. - * - * See LICENSE.txt for license information - ************************************************************************/ - -#include - -#include "TestBed.hpp" - -extern "C" bool mscclUnitTestMode() { return true; } - -namespace RcclUnitTesting -{ - TEST(AllReduce, MscclSingleCall) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllReduce}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclFloat32}; - std::vector const redOps = {ncclSum, ncclProd}; - std::vector const roots = {0}; - std::vector const numElements = {384 * 32 * 32, 384 * 32, 384}; - std::vector const inPlaceList = {true, false}; - std::vector const managedMemList = {true, false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } - - TEST(AllReduce, MscclGroupCall) - { - TestBed testBed; - - // Configuration - ncclFunc_t const funcType = ncclCollAllReduce; - std::vector const& dataTypes = {ncclFloat}; - std::vector const& redOps = {ncclSum}; - std::vector const numElements = {384}; - bool const inPlace = false; - bool const useManagedMem = false; - int const numCollPerGroup = numElements.size(); - - OptionalColArgs options; - // This tests runs 3 collectives in the same group call - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1 && isCorrect; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - // Test either single process all GPUs, or 1 process per GPU - int const numProcesses = isMultiProcess ? totalRanks : 1; - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), numCollPerGroup); - - for (int redOpIdx = 0; redOpIdx < redOps.size() && isCorrect; ++redOpIdx) - { - options.redOp = redOps[redOpIdx]; - for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) - { - if (testBed.ev.showNames) - INFO("%s %d-ranks AllReduce %d Grouped Calls (%s-%s)\n", - isMultiProcess ? "MP" : "SP", - totalRanks, numCollPerGroup, - ncclRedOpNames[redOps[redOpIdx]], ncclDataTypeNames[dataTypes[dataIdx]]); - - // Run all element sizes in parallel as single group - for (int collIdx = 0; collIdx < numCollPerGroup; ++collIdx) - { - testBed.SetCollectiveArgs(funcType, - dataTypes[dataIdx], - numElements[collIdx], - numElements[collIdx], - options, - collIdx); - } - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - testBed.DeallocateMem(); - } - } - testBed.DestroyComms(); - } - testBed.Finalize(); - } - - TEST(AllReduce, MscclPreMultScalar) - { - TestBed testBed; - - // Configuration - ncclFunc_t const funcType = ncclCollAllReduce; - std::vector const& dataTypes = {ncclInt32, ncclFloat32, ncclFloat64}; - ncclRedOp_t const redOp = ncclSum; - std::vector const numElements = {384 * 32 * 32, 384 * 32, 384}; - bool const inPlace = false; - bool const useManagedMem = false; - - OptionalColArgs options; - // Terminate the test as soon as first failure occurs - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - int const numProcesses = isMultiProcess ? totalRanks : 1; - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks)); - - for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) - { - ncclDataType_t const dataType = dataTypes[dataIdx]; - - // Set scalars per rank - PtrUnion scalarsPerRank; - scalarsPerRank.AllocateCpuMem(totalRanks * DataTypeToBytes(dataType)); - for (int i = 0; i < totalRanks; i++) - { - double F = i; - scalarsPerRank.Set(dataType, i, i, F); - } - int const numBytes = totalRanks * DataTypeToBytes(dataType); - memcpy(options.scalarTransport.ptr, scalarsPerRank.ptr, numBytes); - - // Test various scalar residence modes - for (int scalarMode = 0; scalarMode <= 1 && isCorrect; ++scalarMode) - { - if (testBed.ev.showNames) - INFO("%s %d-ranks AllReduce (custom-scalar Mode %d %s)\n", - isMultiProcess ? "MP" : "SP", - totalRanks, scalarMode, ncclDataTypeNames[dataType]); - - for (int i = 0; i < numElements.size() && isCorrect; ++i) - { - options.scalarMode = scalarMode; - options.redOp = redOp; - testBed.SetCollectiveArgs(funcType, dataType, - numElements[i], numElements[i], - options); - // For performance, only allocate and prepare data on largest size - if (i == 0) - { - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - } - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - } - testBed.DeallocateMem(); - } - } - testBed.DestroyComms(); - } - testBed.Finalize(); - } - - TEST(AllReduce, MscclMultiStream) - { - TestBed testBed; - - // Configuration - int const numElements = 384 * 1024; - bool const inPlace = false; - bool const useManagedMem = false; - - OptionalColArgs options; - - // This test runs multiple AllReduce collectives on different streams within the same group call - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1 && isCorrect; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - // Test either single process all GPUs, or 1 process per GPU - int const numProcesses = isMultiProcess ? totalRanks : 1; - - for (int numCollPerGroup = 2; numCollPerGroup <= 6; numCollPerGroup += 2) - { - for (int numStreamsPerGroup = numCollPerGroup; numStreamsPerGroup >= 2; numStreamsPerGroup -= 3) - { - if (testBed.ev.showNames) - INFO("%s %d-ranks Multistream %d-Group Calls across %d streams\n", - isMultiProcess ? "MP" : "SP", totalRanks, numCollPerGroup, numStreamsPerGroup); - - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), - numCollPerGroup, false, numStreamsPerGroup); - - // Set up each collective in group in different stream (modulo numStreamsPerGroup) - options.redOp = ncclSum; - for (int collIdx = 0; collIdx < numCollPerGroup; ++collIdx) - { - testBed.SetCollectiveArgs(ncclCollAllReduce, ncclFloat, numElements, numElements, - options, collIdx, -1, collIdx % numStreamsPerGroup); - } - - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - testBed.DeallocateMem(); - testBed.DestroyComms(); - } - } - } - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllReduce_NonBlockingConf.cpp b/projects/rccl/test/AllReduce_NonBlockingConf.cpp deleted file mode 100644 index b3b8935904..0000000000 --- a/projects/rccl/test/AllReduce_NonBlockingConf.cpp +++ /dev/null @@ -1,64 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" -namespace RcclUnitTesting -{ - TEST(AllReduce, NonBlocking) - { - TestBed testBed; - // Configuration - ncclFunc_t const funcType = ncclCollAllReduce; - std::vector const& dataTypes = {ncclFloat}; - std::vector const& redOps = {ncclSum}; - std::vector const numElements = {1048576}; - bool const inPlace = false; - bool const useManagedMem = false; - bool const useBlocking = false; - - OptionalColArgs options; - // Terminate the test as soon as first failure occurs - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1 && isCorrect; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - // Test either single process all GPUs, or 1 process per GPU - int const numProcesses = isMultiProcess ? totalRanks : 1; - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), 1, useBlocking); - - for (int redOpIdx = 0; redOpIdx < redOps.size() && isCorrect; ++redOpIdx) - { - options.redOp = redOps[redOpIdx]; - for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) - { - if (testBed.ev.showNames) - INFO("%s %d-ranks AllReduce %s Blocking Config (%s-%s)\n", - isMultiProcess ? "MP" : "SP", - totalRanks, useBlocking ? "true" : "false", - ncclRedOpNames[redOps[redOpIdx]], ncclDataTypeNames[dataTypes[dataIdx]]); - - - for (int numIdx = 0; numIdx < numElements.size() && isCorrect; ++numIdx) - { - testBed.SetCollectiveArgs(funcType, - dataTypes[dataIdx], - numElements[numIdx], - numElements[numIdx], - options); - } - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - testBed.DeallocateMem(); - } - } - testBed.DestroyComms(); - } - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllReduce_OutOfPlace.cpp b/projects/rccl/test/AllReduce_OutOfPlace.cpp deleted file mode 100644 index 8bfaf6d39f..0000000000 --- a/projects/rccl/test/AllReduce_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllReduce, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllReduce}; - std::vector const dataTypes = {ncclFloat16, ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclSum, ncclMin}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllReduce_PreMultScalar.cpp b/projects/rccl/test/AllReduce_PreMultScalar.cpp deleted file mode 100644 index 9be19000ad..0000000000 --- a/projects/rccl/test/AllReduce_PreMultScalar.cpp +++ /dev/null @@ -1,80 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - // This tests using custom pre-mult scalars reductions - TEST(AllReduce, PreMultScalar) - { - TestBed testBed; - - // Configuration - ncclFunc_t const funcType = ncclCollAllReduce; - std::vector const& dataTypes = {ncclInt32, ncclFloat32, ncclFloat64}; - ncclRedOp_t const redOp = ncclSum; - std::vector const numElements = {1048576, 1024}; - bool const inPlace = false; - bool const useManagedMem = false; - - OptionalColArgs options; - // Terminate the test as soon as first failure occurs - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - int const numProcesses = isMultiProcess ? totalRanks : 1; - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks)); - - for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) - { - ncclDataType_t const dataType = dataTypes[dataIdx]; - - // Set scalars per rank - PtrUnion scalarsPerRank; - scalarsPerRank.AllocateCpuMem(totalRanks * DataTypeToBytes(dataType)); - for (int i = 0; i < totalRanks; i++) - { - double F = i; - scalarsPerRank.Set(dataType, i, i, F); - } - int const numBytes = totalRanks * DataTypeToBytes(dataType); - memcpy(options.scalarTransport.ptr, scalarsPerRank.ptr, numBytes); - - // Test various scalar residence modes - for (int scalarMode = 0; scalarMode <= 1 && isCorrect; ++scalarMode) - { - if (testBed.ev.showNames) - INFO("%s %d-ranks AllReduce (custom-scalar Mode %d %s)\n", - isMultiProcess ? "MP" : "SP", - totalRanks, scalarMode, ncclDataTypeNames[dataType]); - - for (int i = 0; i < numElements.size() && isCorrect; ++i) - { - options.scalarMode = scalarMode; - options.redOp = redOp; - testBed.SetCollectiveArgs(funcType, dataType, - numElements[i], numElements[i], - options); - // For performance, only allocate and prepare data on largest size - if (i == 0) - { - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - } - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - } - testBed.DeallocateMem(); - } - } - testBed.DestroyComms(); - } - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllToAllTests.cpp b/projects/rccl/test/AllToAllTests.cpp new file mode 100644 index 0000000000..418edb3189 --- /dev/null +++ b/projects/rccl/test/AllToAllTests.cpp @@ -0,0 +1,85 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(AllToAll, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllToAll}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1048576, 1024}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllToAll, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllToAll}; + std::vector const dataTypes = {ncclFloat64, ncclBfloat16}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {5685}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllToAll, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllToAll}; + std::vector const dataTypes = {ncclUint8}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {384 * 1024, 1024}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(AllToAll, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllToAll}; + std::vector const dataTypes = {ncclUint32, ncclUint64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1048576}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/AllToAllVTests.cpp b/projects/rccl/test/AllToAllVTests.cpp new file mode 100644 index 0000000000..bc6fd29706 --- /dev/null +++ b/projects/rccl/test/AllToAllVTests.cpp @@ -0,0 +1,166 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + // Prepare sendcount/recvcounts, sdispls/rdispls arrays within options + void PrepareCounts(int const totalRanks, int const chunkSize, + OptionalColArgs& options, + std::vector& numInputElements, + std::vector& numOutputElements) + { + numInputElements.clear(); + numOutputElements.clear(); + numInputElements.resize(totalRanks, 0); + numOutputElements.resize(totalRanks, 0); + + // Decide how many elements each pair send/recv + for (int sendRank = 0; sendRank < totalRanks; ++sendRank) + for (int recvRank = 0; recvRank < totalRanks; ++recvRank) + { + // Get linear indices into sendcounts/recvcounts array + int const sendIdx = sendRank * totalRanks + recvRank; + int const recvIdx = recvRank * totalRanks + sendRank; + + // Each pair sends slightly different amounts of elements (based on chunkSize) + int const numElements = (1 + sendRank + recvRank) * chunkSize; + options.sendcounts[sendIdx] = options.recvcounts[recvIdx] = numElements; + } + + // Compute displacements + for (int sendRank = 0; sendRank < totalRanks; ++sendRank) + { + int totalSend = 0; + int totalRecv = 0; + + for (int recvRank = 0; recvRank < totalRanks; ++recvRank) + { + int const pairIdx = sendRank * totalRanks + recvRank; + + options.sdispls[pairIdx] = totalSend; + options.rdispls[pairIdx] = totalRecv; + + totalSend += options.sendcounts[pairIdx]; + totalRecv += options.recvcounts[pairIdx]; + } + + numInputElements[sendRank] = totalSend; + numOutputElements[sendRank] = totalRecv; + } + } + + TEST(AllToAllv, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const& dataTypes = {ncclInt32, ncclFloat64, ncclFloat16}; + bool const inPlace = false; + bool const useManagedMem = false; + bool const useHipGraph = false; + + OptionalColArgs options; + + bool isCorrect = true; + for (int totalRanks : testBed.ev.GetNumGpusList()) + for (int isMultiProcess : testBed.ev.GetIsMultiProcessList()) + { + int const numProcesses = isMultiProcess ? totalRanks : 1; + testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks)); + + // Prepare AllToAllV options + std::vector numInputElements; + std::vector numOutputElements; + PrepareCounts(totalRanks, 256, options, numInputElements, numOutputElements); + + for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) + { + if (testBed.ev.showNames) + { + std::string name = testBed.GetTestCaseName(totalRanks, isMultiProcess, + ncclCollAllToAllv, dataTypes[dataIdx], + ncclSum, -1, inPlace, useManagedMem, useHipGraph); + INFO("%s\n", name.c_str()); + } + + for (int rank = 0; rank < totalRanks; ++rank) + { + testBed.SetCollectiveArgs(ncclCollAllToAllv, + dataTypes[dataIdx], + numInputElements[rank], + numOutputElements[rank], + options, + -1, + rank); + } + testBed.AllocateMem(inPlace, useManagedMem); + testBed.PrepareData(); + testBed.ExecuteCollectives({}, useHipGraph); + testBed.ValidateResults(isCorrect); + testBed.DeallocateMem(); + } + testBed.DestroyComms(); + } + testBed.Finalize(); + } + + + TEST(AllToAllv, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const& dataTypes = {ncclFloat32, ncclInt8}; + bool const inPlace = false; + bool const useManagedMem = false; + bool const useHipGraph = false; + + OptionalColArgs options; + + bool isCorrect = true; + for (int totalRanks : testBed.ev.GetNumGpusList()) + for (int isMultiProcess : testBed.ev.GetIsMultiProcessList()) + { + int const numProcesses = isMultiProcess ? totalRanks : 1; + testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks)); + + // Prepare AllToAllV options + std::vector numInputElements; + std::vector numOutputElements; + PrepareCounts(totalRanks, 256, options, numInputElements, numOutputElements); + + for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) + { + if (testBed.ev.showNames) + { + std::string name = testBed.GetTestCaseName(totalRanks, isMultiProcess, + ncclCollAllToAllv, dataTypes[dataIdx], + ncclSum, -1, inPlace, useManagedMem, useHipGraph); + INFO("%s\n", name.c_str()); + } + + for (int rank = 0; rank < totalRanks; ++rank) + { + testBed.SetCollectiveArgs(ncclCollAllToAllv, + dataTypes[dataIdx], + numInputElements[rank], + numOutputElements[rank], + options, + -1, + rank); + } + testBed.AllocateMem(inPlace, useManagedMem); + testBed.PrepareData(); + testBed.ExecuteCollectives({}, useHipGraph); + testBed.ValidateResults(isCorrect); + testBed.DeallocateMem(); + } + testBed.DestroyComms(); + } + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/AllToAll_ManagedMem.cpp b/projects/rccl/test/AllToAll_ManagedMem.cpp deleted file mode 100644 index eecfcc28f7..0000000000 --- a/projects/rccl/test/AllToAll_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllToAll, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllToAll}; - std::vector const dataTypes = {ncclUint8, ncclUint32, ncclUint64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllToAll_OutOfPlace.cpp b/projects/rccl/test/AllToAll_OutOfPlace.cpp deleted file mode 100644 index 4fdd55d6de..0000000000 --- a/projects/rccl/test/AllToAll_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(AllToAll, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollAllToAll}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/AllToAllv_OutOfPlace.cpp b/projects/rccl/test/AllToAllv_OutOfPlace.cpp deleted file mode 100644 index ba4939aacb..0000000000 --- a/projects/rccl/test/AllToAllv_OutOfPlace.cpp +++ /dev/null @@ -1,101 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - - void sendRecvPrep(size_t numInputElementsArray[], - size_t numOutputElementsArray[], - OptionalColArgs &options, - int totalRanks, int numElements) - { - int sendcount = (numElements/totalRanks)*totalRanks; - size_t chunksize = sendcount*2/(totalRanks*totalRanks); - for (int sendRank = 0; sendRank < totalRanks; ++sendRank) - for (int recvRank = 0; recvRank < totalRanks; ++recvRank ) - { - //create send counts, and build other arrays from that - size_t scount = ((sendRank+recvRank)%totalRanks)*chunksize; - if ((sendRank+recvRank)%totalRanks == 0) - scount += (sendcount-chunksize*(totalRanks-1)*totalRanks/2); - options.sendcounts[sendRank*totalRanks+recvRank] = scount; - options.recvcounts[recvRank*totalRanks+sendRank] = options.sendcounts[sendRank*totalRanks+recvRank ]; - } - - for (int sendRank = 0; sendRank < totalRanks; ++sendRank) - { - options.sdispls[sendRank*totalRanks] = 0; - options.rdispls[sendRank*totalRanks] = 0; - for (int recvRank = 1; recvRank < totalRanks; ++recvRank ) - { - options.sdispls[sendRank*totalRanks+recvRank] = - options.sdispls[sendRank*totalRanks+recvRank-1] + options.sendcounts[sendRank*totalRanks+recvRank-1]; - options.rdispls[sendRank*totalRanks+recvRank] = - options.rdispls[sendRank*totalRanks+recvRank-1] + options.recvcounts[sendRank*totalRanks+recvRank-1]; - } - numInputElementsArray[sendRank] = options.sdispls[(sendRank+1)*totalRanks-1] + options.sendcounts[(sendRank+1)*totalRanks-1]; - numOutputElementsArray[sendRank] = options.rdispls[(sendRank+1)*totalRanks-1] + options.recvcounts[(sendRank+1)*totalRanks-1]; - } - - } - - TEST(AllToAllv, OutOfPlace) - { - TestBed testBed; - // Configuration - std::vector const& dataTypes = {ncclInt32, ncclFloat64, ncclFloat16}; - std::vector const numElements = {1048576, 53327, 1024}; - bool const inPlace = false; - bool const useManagedMem = false; - - OptionalColArgs options; - size_t numInputElementsArray[MAX_RANKS], numOutputElementsArray[MAX_RANKS]; - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1 && isCorrect; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - int const numProcesses = isMultiProcess ? totalRanks : 1; - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks)); - - for (int dataIdx = 0; dataIdx < dataTypes.size() && isCorrect; ++dataIdx) - for (int numIdx = 0; numIdx < numElements.size() && isCorrect; ++numIdx) - { - if (testBed.ev.showNames) - { - std::string name = testBed.GetTestCaseName(totalRanks, isMultiProcess, - ncclCollAllToAllv, dataTypes[dataIdx], - ncclSum, -1, - inPlace, useManagedMem, false); - INFO("%s\n", name.c_str()); - - } - sendRecvPrep(numInputElementsArray, numOutputElementsArray, options, totalRanks, numElements[numIdx]); - for (int rank = 0; rank < totalRanks; ++rank) - { - testBed.SetCollectiveArgs(ncclCollAllToAllv, - dataTypes[dataIdx], - numInputElementsArray[rank], - numOutputElementsArray[rank], - options, - -1, - rank); - - } - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - testBed.DeallocateMem(); - - } - testBed.DestroyComms(); - } - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/BroadcastTests.cpp b/projects/rccl/test/BroadcastTests.cpp new file mode 100644 index 0000000000..a50c879c9e --- /dev/null +++ b/projects/rccl/test/BroadcastTests.cpp @@ -0,0 +1,123 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(Broadcast, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1048576, 500}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Broadcast, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast}; + std::vector const dataTypes = {ncclBfloat16, ncclFloat64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {586}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Broadcast, InPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast}; + std::vector const dataTypes = {ncclInt32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {104857, 264}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Broadcast, InPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast}; + std::vector const dataTypes = {ncclInt8, ncclInt64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {958}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Broadcast, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast}; + std::vector const dataTypes = {ncclUint8}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1039203, 2500}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Broadcast, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast}; + std::vector const dataTypes = {ncclUint32, ncclUint64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {896}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/Broadcast_InPlace.cpp b/projects/rccl/test/Broadcast_InPlace.cpp deleted file mode 100644 index 694c3c84d4..0000000000 --- a/projects/rccl/test/Broadcast_InPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Broadcast, InPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollBroadcast}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Broadcast_ManagedMem.cpp b/projects/rccl/test/Broadcast_ManagedMem.cpp deleted file mode 100644 index 66f5122e49..0000000000 --- a/projects/rccl/test/Broadcast_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Broadcast, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollBroadcast}; - std::vector const dataTypes = {ncclUint8, ncclUint32, ncclUint64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Broadcast_OutOfPlace.cpp b/projects/rccl/test/Broadcast_OutOfPlace.cpp deleted file mode 100644 index e52f429370..0000000000 --- a/projects/rccl/test/Broadcast_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Broadcast, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollBroadcast}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {1}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index b405660d9d..0b00f33cea 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -44,60 +44,22 @@ if(BUILD_TESTS) # Collect source files for tests if(BUILD_ALLREDUCE_ONLY) set(TEST_SOURCE_FILES - AllReduce_Clique.cpp - AllReduce_GroupCall.cpp - AllReduce_NonBlockingConf.cpp - AllReduce_InPlace.cpp - AllReduce_ManagedMem.cpp - AllReduce_OutOfPlace.cpp - AllReduce_PreMultScalar.cpp - AllReduce_Msccl.cpp - Multistream.cpp + AllReduce_Tests.cpp ) else() set(TEST_SOURCE_FILES - #AllReduce - AllReduce_Clique.cpp - AllReduce_GroupCall.cpp - AllReduce_NonBlockingConf.cpp - AllReduce_InPlace.cpp - AllReduce_ManagedMem.cpp - AllReduce_OutOfPlace.cpp - AllReduce_PreMultScalar.cpp - AllReduce_Msccl.cpp - #AllGather - AllGather_InPlace.cpp - AllGather_ManagedMem.cpp - AllGather_OutOfPlace.cpp - #AllToAll - AllToAll_OutOfPlace.cpp - AllToAll_ManagedMem.cpp - #AllToAllv - AllToAllv_OutOfPlace.cpp - #Broadcast - Broadcast_InPlace.cpp - Broadcast_ManagedMem.cpp - Broadcast_OutOfPlace.cpp - #Reduce - Reduce_InPlace.cpp - Reduce_ManagedMem.cpp - Reduce_OutOfPlace.cpp - #ReduceScatter - ReduceScatter_InPlace.cpp - ReduceScatter_ManagedMem.cpp - ReduceScatter_OutOfPlace.cpp - #Scatter - Scatter_InPlace.cpp - Scatter_ManagedMem.cpp - Scatter_OutOfPlace.cpp - #Gather - Gather_InPlace.cpp - Gather_ManagedMem.cpp - Gather_OutOfPlace.cpp - #SendRecv - SendRecv_SinglePairs.cpp - #Misc - Multistream.cpp + AllGatherTests.cpp + AllReduceTests.cpp + AllToAllTests.cpp + AllToAllVTests.cpp + BroadcastTests.cpp + GatherTests.cpp + GroupCallTests.cpp + NonBlockingTests.cpp + ReduceScatterTests.cpp + ReduceTests.cpp + ScatterTests.cpp + SendRecvTests.cpp ) endif() diff --git a/projects/rccl/test/GatherTests.cpp b/projects/rccl/test/GatherTests.cpp new file mode 100644 index 0000000000..17aad21be9 --- /dev/null +++ b/projects/rccl/test/GatherTests.cpp @@ -0,0 +1,124 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(Gather, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollGather}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {1048576, 127}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Gather, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollGather}; + std::vector const dataTypes = {ncclFloat16, ncclFloat64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {924873}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + + TEST(Gather, InPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollGather}; + std::vector const dataTypes = {ncclInt8, ncclInt32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {13576}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Gather, InPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollGather}; + std::vector const dataTypes = {ncclInt64, ncclFloat16}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1048576}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Gather, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollGather}; + std::vector const dataTypes = {ncclUint8, ncclUint32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {1051234}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Gather, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollGather}; + std::vector const dataTypes = {ncclUint64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {5231}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/Gather_InPlace.cpp b/projects/rccl/test/Gather_InPlace.cpp deleted file mode 100644 index a2accce3e1..0000000000 --- a/projects/rccl/test/Gather_InPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Gather, InPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollGather}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64, ncclFloat16}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Gather_ManagedMem.cpp b/projects/rccl/test/Gather_ManagedMem.cpp deleted file mode 100644 index 837c863544..0000000000 --- a/projects/rccl/test/Gather_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Gather, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollGather}; - std::vector const dataTypes = {ncclUint8, ncclUint32, ncclUint64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {1}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Gather_OutOfPlace.cpp b/projects/rccl/test/Gather_OutOfPlace.cpp deleted file mode 100644 index 6a018ec160..0000000000 --- a/projects/rccl/test/Gather_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Gather, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollGather}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {1}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/GroupCallTests.cpp b/projects/rccl/test/GroupCallTests.cpp new file mode 100644 index 0000000000..52ca200203 --- /dev/null +++ b/projects/rccl/test/GroupCallTests.cpp @@ -0,0 +1,170 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + // Test identical collectives within the same group call + TEST(GroupCall, Identical) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollAllReduce, ncclCollAllReduce, ncclCollAllReduce}; + std::vector const redOps = {ncclSum, ncclSum, ncclSum}; + std::vector const dataTypes = {ncclFloat, ncclFloat, ncclFloat}; + std::vector const numElements = {1048576, 384 * 1024, 384}; + + int const numCollPerGroup = numElements.size(); + bool const inPlace = false; + bool const useManagedMem = false; + + bool isCorrect = true; + for (int totalRanks : testBed.ev.GetNumGpusList()) + for (int isMultiProcess : testBed.ev.GetIsMultiProcessList()) + { + // Test either single process all GPUs, or 1 process per GPU + int const numProcesses = isMultiProcess ? totalRanks : 1; + testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), numCollPerGroup); + + if (testBed.ev.showNames) + INFO("%s %d-ranks GroupCall Identical\n", isMultiProcess ? "MP" : "SP", totalRanks); + + // Set up the different collectives within the group + for (int collIdx = 0; collIdx < numCollPerGroup; ++collIdx) + { + OptionalColArgs options; + options.redOp = redOps[collIdx]; + testBed.SetCollectiveArgs(funcTypes[collIdx], + dataTypes[collIdx], + numElements[collIdx], + numElements[collIdx], + options, + collIdx); + } + + testBed.AllocateMem(inPlace, useManagedMem); + testBed.PrepareData(); + testBed.ExecuteCollectives(); + testBed.ValidateResults(isCorrect); + testBed.DeallocateMem(); + testBed.DestroyComms(); + } + testBed.Finalize(); + } + + // Test different collectives within the same group call + TEST(GroupCall, Different) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast, + ncclCollAllGather, + ncclCollReduceScatter, + ncclCollAllReduce, + ncclCollGather, + ncclCollScatter, + ncclCollAllToAll}; + int const numCollPerGroup = funcTypes.size(); + int const numElements = 1048576; + bool const inPlace = false; + bool const useManagedMem = false; + + OptionalColArgs options; + options.redOp = ncclSum; + options.root = 0; + + bool isCorrect = true; + for (int totalRanks : testBed.ev.GetNumGpusList()) + for (int isMultiProcess : testBed.ev.GetIsMultiProcessList()) + { + // Test either single process all GPUs, or 1 process per GPU + int const numProcesses = isMultiProcess ? totalRanks : 1; + testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), numCollPerGroup); + + if (testBed.ev.showNames) + INFO("%s %d-ranks GroupCall Different\n", isMultiProcess ? "MP" : "SP", totalRanks); + + // Set up the different collectives within the group + for (int collIdx = 0; collIdx < numCollPerGroup; ++collIdx) + { + int numInputElements; + int numOutputElements; + CollectiveArgs::GetNumElementsForFuncType(funcTypes[collIdx], + numElements, + totalRanks, + &numInputElements, + &numOutputElements); + + testBed.SetCollectiveArgs(funcTypes[collIdx], + ncclFloat, + numInputElements, + numOutputElements, + options, + collIdx); + } + + testBed.AllocateMem(inPlace, useManagedMem); + testBed.PrepareData(); + testBed.ExecuteCollectives(); + testBed.ValidateResults(isCorrect); + testBed.DeallocateMem(); + testBed.DestroyComms(); + } + testBed.Finalize(); + } + + TEST(GroupCall, Multistream) + { + TestBed testBed; + + // Configuration + int const numElements = 1048576; + bool const inPlace = false; + bool const useManagedMem = false; + + OptionalColArgs options; + + // This test runs multiple AllReduce collectives on different streams within the same group call + bool isCorrect = true; + for (int totalRanks : testBed.ev.GetNumGpusList()) + for (int isMultiProcess : testBed.ev.GetIsMultiProcessList()) + { + // Test either single process all GPUs, or 1 process per GPU + int const numProcesses = isMultiProcess ? totalRanks : 1; + + for (int numCollPerGroup = 2; numCollPerGroup <= 6; numCollPerGroup += 2) + { + for (int numStreamsPerGroup = numCollPerGroup; numStreamsPerGroup >= 2; numStreamsPerGroup -= 3) + { + if (testBed.ev.showNames) + INFO("%s %d-ranks Multistream %d-Group Calls across %d streams\n", + isMultiProcess ? "MP" : "SP", totalRanks, numCollPerGroup, numStreamsPerGroup); + + testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), + numCollPerGroup, true, numStreamsPerGroup); + + // Set up each collective in group in different stream (modulo numStreamsPerGroup) + options.redOp = ncclSum; + for (int collIdx = 0; collIdx < numCollPerGroup; ++collIdx) + { + testBed.SetCollectiveArgs(ncclCollAllReduce, ncclFloat, numElements, numElements, + options, collIdx, -1, collIdx % numStreamsPerGroup); + } + + testBed.AllocateMem(inPlace, useManagedMem); + testBed.PrepareData(); + testBed.ExecuteCollectives(); + testBed.ValidateResults(isCorrect); + testBed.DeallocateMem(); + testBed.DestroyComms(); + } + } + } + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/Multistream.cpp b/projects/rccl/test/Multistream.cpp deleted file mode 100644 index fa601ccbfb..0000000000 --- a/projects/rccl/test/Multistream.cpp +++ /dev/null @@ -1,61 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" -#include -namespace RcclUnitTesting -{ - TEST(Multistream, NoGraph) - { - TestBed testBed; - - // Configuration - int const numElements = 1048576; - bool const inPlace = false; - bool const useManagedMem = false; - - OptionalColArgs options; - - // This test runs multiple AllReduce collectives on different streams within the same group call - bool isCorrect = true; - for (int totalRanks = testBed.ev.minGpus; totalRanks <= testBed.ev.maxGpus && isCorrect; ++totalRanks) - for (int isMultiProcess = 0; isMultiProcess <= 1 && isCorrect; ++isMultiProcess) - { - if (!(testBed.ev.processMask & (1 << isMultiProcess))) continue; - - // Test either single process all GPUs, or 1 process per GPU - int const numProcesses = isMultiProcess ? totalRanks : 1; - - for (int numCollPerGroup = 2; numCollPerGroup <= 6; numCollPerGroup += 2) - { - for (int numStreamsPerGroup = numCollPerGroup; numStreamsPerGroup >= 2; numStreamsPerGroup -= 3) - { - if (testBed.ev.showNames) - INFO("%s %d-ranks Multistream %d-Group Calls across %d streams\n", - isMultiProcess ? "MP" : "SP", totalRanks, numCollPerGroup, numStreamsPerGroup); - - testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), - numCollPerGroup, false, numStreamsPerGroup); - - // Set up each collective in group in different stream (modulo numStreamsPerGroup) - options.redOp = ncclSum; - for (int collIdx = 0; collIdx < numCollPerGroup; ++collIdx) - { - testBed.SetCollectiveArgs(ncclCollAllReduce, ncclFloat, numElements, numElements, - options, collIdx, -1, collIdx % numStreamsPerGroup); - } - - testBed.AllocateMem(inPlace, useManagedMem); - testBed.PrepareData(); - testBed.ExecuteCollectives(); - testBed.ValidateResults(isCorrect); - testBed.DeallocateMem(); - testBed.DestroyComms(); - } - } - } - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/NonBlockingTests.cpp b/projects/rccl/test/NonBlockingTests.cpp new file mode 100644 index 0000000000..fab505ad70 --- /dev/null +++ b/projects/rccl/test/NonBlockingTests.cpp @@ -0,0 +1,70 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + // Test various collectives using with non-blocking comms + TEST(NonBlocking, SingleCalls) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollBroadcast, + ncclCollReduce, + ncclCollAllGather, + ncclCollReduceScatter, + ncclCollAllReduce, + ncclCollGather, + ncclCollScatter}; + int const numElements = 1048576; + bool const inPlace = false; + bool const useManagedMem = false; + bool const useBlocking = false; + + OptionalColArgs options; + options.redOp = ncclSum; + + bool isCorrect = true; + for (int totalRanks : testBed.ev.GetNumGpusList()) + for (int isMultiProcess : testBed.ev.GetIsMultiProcessList()) + { + int const numProcesses = isMultiProcess ? totalRanks : 1; + // Initialize communicators in non-blocking mode + testBed.InitComms(TestBed::GetDeviceIdsList(numProcesses, totalRanks), 1, useBlocking); + + // Loop over various collective functions + for (auto funcType : funcTypes) + { + if (testBed.ev.showNames) + INFO("%s %d-ranks Non-Blocking %s\n", + isMultiProcess ? "MP" : "SP", totalRanks, ncclFuncNames[funcType]); + + int numInputElements; + int numOutputElements; + CollectiveArgs::GetNumElementsForFuncType(funcType, + numElements, + totalRanks, + &numInputElements, + &numOutputElements); + + testBed.SetCollectiveArgs(funcType, + ncclFloat, + numInputElements, + numOutputElements, + options); + + testBed.AllocateMem(inPlace, useManagedMem); + testBed.PrepareData(); + testBed.ExecuteCollectives(); + testBed.ValidateResults(isCorrect); + testBed.DeallocateMem(); + } + testBed.DestroyComms(); + } + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/ReduceScatterTests.cpp b/projects/rccl/test/ReduceScatterTests.cpp new file mode 100644 index 0000000000..ef7f8f701b --- /dev/null +++ b/projects/rccl/test/ReduceScatterTests.cpp @@ -0,0 +1,123 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(ReduceScatter, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduceScatter}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclMax}; + std::vector const roots = {0}; + std::vector const numElements = {393216, 384}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(ReduceScatter, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduceScatter}; + std::vector const dataTypes = {ncclFloat64, ncclBfloat16}; + std::vector const redOps = {ncclMax}; + std::vector const roots = {0}; + std::vector const numElements = {1048576}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(ReduceScatter, InPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduceScatter}; + std::vector const dataTypes = {ncclInt32}; + std::vector const redOps = {ncclProd}; + std::vector const roots = {0, 1}; + std::vector const numElements = {542357}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(ReduceScatter, InPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduceScatter}; + std::vector const dataTypes = {ncclUint8, ncclFloat16}; + std::vector const redOps = {ncclMin}; + std::vector const roots = {0}; + std::vector const numElements = {246};; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(ReduceScatter, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduceScatter}; + std::vector const dataTypes = {ncclInt64, ncclUint8}; + std::vector const redOps = {ncclAvg}; + std::vector const roots = {0}; + std::vector const numElements = {1024}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(ReduceScatter, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduceScatter}; + std::vector const dataTypes = {ncclUint32, ncclUint64}; + std::vector const redOps = {ncclAvg}; + std::vector const roots = {0}; + std::vector const numElements = {6485423}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/ReduceScatter_InPlace.cpp b/projects/rccl/test/ReduceScatter_InPlace.cpp deleted file mode 100644 index 649b2bf524..0000000000 --- a/projects/rccl/test/ReduceScatter_InPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(ReduceScatter, InPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollReduceScatter}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64}; - std::vector const redOps = {ncclProd}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024};; - std::vector const inPlaceList = {true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/ReduceScatter_ManagedMem.cpp b/projects/rccl/test/ReduceScatter_ManagedMem.cpp deleted file mode 100644 index 87659a98ee..0000000000 --- a/projects/rccl/test/ReduceScatter_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(ReduceScatter, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollReduceScatter}; - std::vector const dataTypes = {ncclUint8, ncclUint32, ncclUint64}; - std::vector const redOps = {ncclAvg}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/ReduceScatter_OutOfPlace.cpp b/projects/rccl/test/ReduceScatter_OutOfPlace.cpp deleted file mode 100644 index 714f7c3db9..0000000000 --- a/projects/rccl/test/ReduceScatter_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(ReduceScatter, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollReduceScatter}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclMax}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/ReduceTests.cpp b/projects/rccl/test/ReduceTests.cpp new file mode 100644 index 0000000000..73280b0952 --- /dev/null +++ b/projects/rccl/test/ReduceTests.cpp @@ -0,0 +1,123 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(Reduce, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduce}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {393216, 384}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Reduce, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduce}; + std::vector const dataTypes = {ncclFloat16, ncclFloat64}; + std::vector const redOps = {ncclMin}; + std::vector const roots = {0}; + std::vector const numElements = {393216}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Reduce, InPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduce}; + std::vector const dataTypes = {ncclInt32, ncclInt8}; + std::vector const redOps = {ncclProd}; + std::vector const roots = {1}; + std::vector const numElements = {384}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Reduce, InPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduce}; + std::vector const dataTypes = {ncclBfloat16}; + std::vector const redOps = {ncclMax}; + std::vector const roots = {0}; + std::vector const numElements = {393216}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Reduce, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduce}; + std::vector const dataTypes = {ncclUint64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {3524082, 2500}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Reduce, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollReduce}; + std::vector const dataTypes = {ncclFloat64, ncclBfloat16}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {4314}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/Reduce_InPlace.cpp b/projects/rccl/test/Reduce_InPlace.cpp deleted file mode 100644 index b2f417c5bb..0000000000 --- a/projects/rccl/test/Reduce_InPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Reduce, InPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollReduce}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64}; - std::vector const redOps = {ncclSum, ncclProd}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Reduce_ManagedMem.cpp b/projects/rccl/test/Reduce_ManagedMem.cpp deleted file mode 100644 index 003b2ca18a..0000000000 --- a/projects/rccl/test/Reduce_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Reduce, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollReduce}; - std::vector const dataTypes = {ncclUint8, ncclUint32, ncclUint64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Reduce_OutOfPlace.cpp b/projects/rccl/test/Reduce_OutOfPlace.cpp deleted file mode 100644 index 0900911f0b..0000000000 --- a/projects/rccl/test/Reduce_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Reduce, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollReduce}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclMin, ncclMax, ncclAvg}; - std::vector const roots = {1}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/ScatterTests.cpp b/projects/rccl/test/ScatterTests.cpp new file mode 100644 index 0000000000..c0b906475b --- /dev/null +++ b/projects/rccl/test/ScatterTests.cpp @@ -0,0 +1,123 @@ +/************************************************************************* + * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ +#include "TestBed.hpp" + +namespace RcclUnitTesting +{ + TEST(Scatter, OutOfPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollScatter}; + std::vector const dataTypes = {ncclFloat32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {393216, 384}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Scatter, OutOfPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollScatter}; + std::vector const dataTypes = {ncclFloat64, ncclBfloat16}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {1}; + std::vector const numElements = {24658}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Scatter, InPlace) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollScatter}; + std::vector const dataTypes = {ncclInt32}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {1048576, 1024}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Scatter, InPlaceGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollScatter}; + std::vector const dataTypes = {ncclInt8, ncclFloat16}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {356}; + std::vector const inPlaceList = {true}; + std::vector const managedMemList = {false}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Scatter, ManagedMem) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollScatter}; + std::vector const dataTypes = {ncclInt64, ncclUint8}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {948576}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {false}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } + + TEST(Scatter, ManagedMemGraph) + { + TestBed testBed; + + // Configuration + std::vector const funcTypes = {ncclCollScatter}; + std::vector const dataTypes = {ncclUint32, ncclUint64}; + std::vector const redOps = {ncclSum}; + std::vector const roots = {0}; + std::vector const numElements = {125}; + std::vector const inPlaceList = {false}; + std::vector const managedMemList = {true}; + std::vector const useHipGraphList = {true}; + + testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, + inPlaceList, managedMemList, useHipGraphList); + testBed.Finalize(); + } +} diff --git a/projects/rccl/test/Scatter_InPlace.cpp b/projects/rccl/test/Scatter_InPlace.cpp deleted file mode 100644 index 2b1eda72b0..0000000000 --- a/projects/rccl/test/Scatter_InPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Scatter, InPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollScatter}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {true}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Scatter_ManagedMem.cpp b/projects/rccl/test/Scatter_ManagedMem.cpp deleted file mode 100644 index 1383a0c96d..0000000000 --- a/projects/rccl/test/Scatter_ManagedMem.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Scatter, ManagedMem) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollScatter}; - std::vector const dataTypes = {ncclUint8, ncclUint32, ncclUint64}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {true}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/Scatter_OutOfPlace.cpp b/projects/rccl/test/Scatter_OutOfPlace.cpp deleted file mode 100644 index f7facf3959..0000000000 --- a/projects/rccl/test/Scatter_OutOfPlace.cpp +++ /dev/null @@ -1,28 +0,0 @@ -/************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - * - * See LICENSE.txt for license information - ************************************************************************/ -#include "TestBed.hpp" - -namespace RcclUnitTesting -{ - TEST(Scatter, OutOfPlace) - { - TestBed testBed; - - // Configuration - std::vector const funcTypes = {ncclCollScatter}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; - std::vector const redOps = {ncclSum}; - std::vector const roots = {1}; - std::vector const numElements = {1048576, 53327, 1024}; - std::vector const inPlaceList = {false}; - std::vector const managedMemList = {false}; - std::vector const useHipGraphList = {false, true}; - - testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements, - inPlaceList, managedMemList, useHipGraphList); - testBed.Finalize(); - } -} diff --git a/projects/rccl/test/SendRecv_SinglePairs.cpp b/projects/rccl/test/SendRecvTests.cpp similarity index 97% rename from projects/rccl/test/SendRecv_SinglePairs.cpp rename to projects/rccl/test/SendRecvTests.cpp index 3a4ad38206..fe59459151 100644 --- a/projects/rccl/test/SendRecv_SinglePairs.cpp +++ b/projects/rccl/test/SendRecvTests.cpp @@ -1,5 +1,5 @@ /************************************************************************* - * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved. * * See LICENSE.txt for license information ************************************************************************/ diff --git a/projects/rccl/test/common/EnvVars.cpp b/projects/rccl/test/common/EnvVars.cpp index 9f783c14ed..c27560f20e 100644 --- a/projects/rccl/test/common/EnvVars.cpp +++ b/projects/rccl/test/common/EnvVars.cpp @@ -56,6 +56,7 @@ namespace RcclUnitTesting showNames = GetEnvVar("UT_SHOW_NAMES" , 1); minGpus = GetEnvVar("UT_MIN_GPUS" , 2); maxGpus = GetEnvVar("UT_MAX_GPUS" , numDevicesAvailable); + onlyPow2Gpus = GetEnvVar("UT_POW2_GPUS" , false); processMask = GetEnvVar("UT_PROCESS_MASK", UT_SINGLE_PROCESS | UT_MULTI_PROCESS); verbose = GetEnvVar("UT_VERBOSE" , 0); printValues = GetEnvVar("UT_PRINT_VALUES", 0); @@ -124,6 +125,17 @@ namespace RcclUnitTesting dataTypes.push_back(ncclBfloat16); #endif } + + // Build list of possible # GPU ranks based on env vars + numGpusList.clear(); + for (int i = minGpus; i <= maxGpus; i++) + if (!onlyPow2Gpus || ((i & (i-1)) == 0)) + numGpusList.push_back(i); + + // Build isMultiProcessList + isMultiProcessList.clear(); + if (this->processMask & UT_SINGLE_PROCESS) isMultiProcessList.push_back(0); + if (this->processMask & UT_MULTI_PROCESS) isMultiProcessList.push_back(1); } std::vector const& EnvVars::GetAllSupportedRedOps() @@ -136,6 +148,16 @@ namespace RcclUnitTesting return dataTypes; } + std::vector const& EnvVars::GetNumGpusList() + { + return numGpusList; + } + + std::vector const& EnvVars::GetIsMultiProcessList() + { + return isMultiProcessList; + } + int EnvVars::GetEnvVar(std::string const varname, int defaultValue) { if (getenv(varname.c_str())) @@ -165,6 +187,7 @@ namespace RcclUnitTesting std::make_pair("UT_SHOW_NAMES" , "Show test case names"), std::make_pair("UT_MIN_GPUS" , "Minimum number of GPUs to use"), std::make_pair("UT_MAX_GPUS" , "Maximum number of GPUs to use"), + std::make_pair("UT_POW2_GPUS" , "Only allow power-of-2 # of GPUs"), std::make_pair("UT_PROCESS_MASK" , "Whether to run single/multi process"), std::make_pair("UT_VERBOSE" , "Show verbose unit test output"), std::make_pair("UT_REDOPS" , "List of reduction ops to test"), diff --git a/projects/rccl/test/common/EnvVars.hpp b/projects/rccl/test/common/EnvVars.hpp index eeef060acc..a3a510e3c0 100644 --- a/projects/rccl/test/common/EnvVars.hpp +++ b/projects/rccl/test/common/EnvVars.hpp @@ -21,6 +21,7 @@ namespace RcclUnitTesting bool showNames; // List test case names during run [UT_SHOW_NAMES] int minGpus; // Set the minimum number of GPUs to use [UT_MIN_GPUS] int maxGpus; // Set the maximum number of GPUs to use [UT_MAX_GPUS] + bool onlyPow2Gpus; // Only allow power-of-2 # of GPUs [UT_POW2_GPUS] int processMask; // Filter single/multi process [UT_PROCESS_MASK] bool verbose; // Show verbose TestBed output for debug [UT_VERBOSE] int printValues; // Print out input/output/expected arrays [UT_PRINT_VALUES] @@ -34,11 +35,15 @@ namespace RcclUnitTesting std::vector const& GetAllSupportedRedOps(); std::vector const& GetAllSupportedDataTypes(); + std::vector const& GetNumGpusList(); + std::vector const& GetIsMultiProcessList(); static void ShowConfig(); protected: - std::vector redOps; // Supported reduction ops [UT_REDOPS] - std::vector dataTypes; // Support datatypes [UT_DATATYPES] + std::vector redOps; // Supported reduction ops [UT_REDOPS] + std::vector dataTypes; // Support datatypes [UT_DATATYPES] + std::vector numGpusList; // List of # Gpus to use [UT_MIN_GPUS/UT_MAX_GPUS/UT_POW2_GPUS] + std::vector isMultiProcessList; // Single or multi process [UT_PROCESS_MASK] // Helper functions to parse environment variables int GetEnvVar(std::string const varname, int defaultValue); diff --git a/projects/rccl/test/common/TestBed.cpp b/projects/rccl/test/common/TestBed.cpp index 92187b39f7..f293d4d6f8 100644 --- a/projects/rccl/test/common/TestBed.cpp +++ b/projects/rccl/test/common/TestBed.cpp @@ -55,34 +55,6 @@ namespace RcclUnitTesting // Collect the number of GPUs this->numDevicesAvailable = ev.maxGpus; if (ev.verbose) INFO("Detected %d GPUs\n", this->numDevicesAvailable); - - // Create the maximum number of possible child processes (1 per GPU) - // Parent and child communicate via pipes - childList.resize(this->numDevicesAvailable); - for (int childId = 0; childId < this->numDevicesAvailable; ++childId) - { - childList[childId] = new TestBedChild(childId, ev.verbose, ev.printValues); - if (childList[childId]->InitPipes() != TEST_SUCCESS) - { - ERROR("Unable to create pipes to child process\n"); - return; - } - - pid_t pid = fork(); - if (pid == 0) - { - // Child process enters execution loop - childList[childId]->StartExecutionLoop(); - return; - } - else - { - // Parent records child process ID and closes unused ends of pipe - childList[childId]->pid = pid; - close(childList[childId]->childWriteFd); - close(childList[childId]->childReadFd); - } - } } void TestBed::InitComms(std::vector> const& deviceIdsPerProcess, @@ -112,6 +84,40 @@ namespace RcclUnitTesting } } + // Check that no children currently exist + if (childList.size() > 0) + { + ERROR("DestroyComms must be called prior to subsequent call to InitComms\n"); + return; + } + + // Create child-processes + childList.resize(this->numDevicesAvailable); + for (int childId = 0; childId < this->numDevicesAvailable; ++childId) + { + childList[childId] = new TestBedChild(childId, ev.verbose, ev.printValues); + if (childList[childId]->InitPipes() != TEST_SUCCESS) + { + ERROR("Unable to create pipes to child process\n"); + return; + } + + pid_t pid = fork(); + if (pid == 0) + { + // Child process enters execution loop + childList[childId]->StartExecutionLoop(); + return; + } + else + { + // Parent records child process ID and closes unused ends of pipe + childList[childId]->pid = pid; + close(childList[childId]->childWriteFd); + close(childList[childId]->childReadFd); + } + } + // Determine number of unique GPUs being used. std::set unique_devices; for (auto a: this->rankToDeviceMap) @@ -375,17 +381,19 @@ namespace RcclUnitTesting PIPE_CHECK(childId); } - // Reset bookkeeping - this->numActiveChildren = 0; - this->numActiveRanks = 0; - this->numCollectivesInGroup = 0; + // Close any open child processes + Finalize(); InteractiveWait("Finishing DestroyComms"); } void TestBed::Finalize() { + if (this->numActiveChildren == 0) + return; + InteractiveWait("Starting Finalize"); + // Send Stop to all child processes int const cmd = TestBedChild::CHILD_STOP; for (int childId = 0; childId < this->numDevicesAvailable; ++childId) @@ -396,7 +404,25 @@ namespace RcclUnitTesting close(childList[childId]->parentWriteFd); close(childList[childId]->parentReadFd); } - this->numDevicesAvailable = 0; + + // Wait for processes to stop + for (int childId = 0; childId < this->numActiveChildren; ++childId) + { + int returnVal = 0; + waitpid(childList[childId]->pid, &returnVal, 0); + if (returnVal != 0) + { + ERROR("Child process %d exited with code %d\n", childId, returnVal); + } + } + + childList.clear(); + + // Reset bookkeeping + this->numActiveChildren = 0; + this->numActiveRanks = 0; + this->numCollectivesInGroup = 0; + InteractiveWait("Finishing Finalize"); } @@ -455,12 +481,12 @@ namespace RcclUnitTesting else ss << " "; ss << "ranks "; - ss << ncclFuncNames[funcType] << " "; + ss << std::setfill(' ') << std::setw(20) << ncclFuncNames[funcType] << " "; ss << "(" << (inPlace ? "IP" : "OP") << "," << (managedMem ? "MM" : "GM") << "," << (useHipGraph ? "GL" : "NL") <<") "; - ss << ncclDataTypeNames[dataType] << " "; - if (CollectiveArgs::UsesReduce(funcType)) ss << ncclRedOpNames[redOp] << " "; + ss << std::setfill(' ') << std::setw(12) << ncclDataTypeNames[dataType] << " "; + if (CollectiveArgs::UsesReduce(funcType)) ss << std::setfill(' ') << std::setw(7) << ncclRedOpNames[redOp] << " "; if (CollectiveArgs::UsesRoot(funcType)) ss << "Root " << root << " "; return ss.str(); } @@ -511,17 +537,19 @@ namespace RcclUnitTesting bool isCorrect = true; // Sweep over the number of ranks - for (int ranksPerGpu=1; ranksPerGpu <= ev.maxRanksPerGpu; ranksPerGpu++) - for (int numGpus = ev.minGpus; numGpus <= ev.maxGpus && isCorrect; ++numGpus) - for (int isMultiProcess = 0; isMultiProcess <= 1 && isCorrect; ++isMultiProcess) + for (int numGpus : ev.GetNumGpusList()) + for (int isMultiProcess : ev.GetIsMultiProcessList()) + for (int ranksPerGpu=1; ranksPerGpu <= ev.maxRanksPerGpu && isCorrect; ++ranksPerGpu) { - if (!(ev.processMask & (1 << isMultiProcess))) continue; - // Test either single process all GPUs, or 1 process per GPU int const numChildren = isMultiProcess ? numGpus : 1; int const numRanks = numGpus*ranksPerGpu; this->InitComms(TestBed::GetDeviceIdsList(numChildren, numGpus, ranksPerGpu)); - if (testing::Test::HasFailure()) continue; + if (testing::Test::HasFailure()) + { + isCorrect = false; + continue; + } for (int ftIdx = 0; ftIdx < funcTypes.size() && isCorrect; ++ftIdx) for (int dtIdx = 0; dtIdx < dataTypes.size() && isCorrect; ++dtIdx) @@ -545,13 +573,21 @@ namespace RcclUnitTesting numInputElements, numOutputElements, optionalArgs); - if (testing::Test::HasFailure()) continue; + if (testing::Test::HasFailure()) + { + isCorrect = false; + continue; + } // Only allocate once for largest size if (neIdx == 0) { this->AllocateMem(inPlaceList[ipIdx], managedMemList[mmIdx]); - if (testing::Test::HasFailure()) continue; + if (testing::Test::HasFailure()) + { + isCorrect = false; + continue; + } } for (int hgIdx = 0; hgIdx < useHipGraphList.size() && isCorrect; ++hgIdx) @@ -563,7 +599,11 @@ namespace RcclUnitTesting funcTypes[ftIdx] == ncclCollReduce || funcTypes[ftIdx] == ncclCollAllReduce)); if (!canSkip) this->PrepareData(); - if (testing::Test::HasFailure()) continue; + if (testing::Test::HasFailure()) + { + isCorrect = false; + continue; + } std::string name = this->GetTestCaseName(numGpus, isMultiProcess, funcTypes[ftIdx], dataTypes[dtIdx], @@ -573,12 +613,16 @@ namespace RcclUnitTesting if (ev.showNames) { - INFO("%s [%d elements]\n", name.c_str(), numInputElements); + INFO("%s [%9d elements]\n", name.c_str(), numInputElements); } std::vector currentRanksEmpty = {}; this->ExecuteCollectives(currentRanksEmpty, useHipGraphList[hgIdx]); - if (testing::Test::HasFailure()) continue; + if (testing::Test::HasFailure()) + { + isCorrect = false; + continue; + } this->ValidateResults(isCorrect); if (!isCorrect) { diff --git a/projects/rccl/test/common/TestBedChild.cpp b/projects/rccl/test/common/TestBedChild.cpp index dc2ed066a3..a45f8bdedb 100644 --- a/projects/rccl/test/common/TestBedChild.cpp +++ b/projects/rccl/test/common/TestBedChild.cpp @@ -99,7 +99,7 @@ namespace RcclUnitTesting case CHILD_VALIDATE_RESULTS: status = ValidateResults(); break; case CHILD_DEALLOCATE_MEM : status = DeallocateMem(); break; case CHILD_DESTROY_COMMS : status = DestroyComms(); break; - case CHILD_STOP : status = Stop(); break; + case CHILD_STOP : goto stop; default: exit(0); } @@ -112,6 +112,7 @@ namespace RcclUnitTesting break; } } + stop: if (verbose) INFO("Child %d exiting execution loop\n", this->childId); // Close child ends of pipe @@ -433,6 +434,7 @@ namespace RcclUnitTesting { CHECK_HIP(hipSetDevice(this->deviceIds[localRank])); if (this->verbose) INFO("Capturing stream for rank %d\n", localRank); + CHECK_HIP(hipSetDevice(this->deviceIds[localRank])); for (int i = 0; i < this->numStreamsPerGroup; i++) { CHECK_HIP(hipStreamBeginCapture(this->streams[localRank][i], hipStreamCaptureModeRelaxed)); @@ -686,7 +688,7 @@ namespace RcclUnitTesting for (int localRank : localRanksToExecute) { CollectiveArgs const& collArg = this->collArgs[localRank][collId]; - + CHECK_HIP(hipSetDevice(this->deviceIds[localRank])); int numOutputElementsToPrint = (this->printValues < 0 ? collArg.numOutputElements : this->printValues); size_t const numOutputBytes = numOutputElementsToPrint * DataTypeToBytes(collArg.dataType); CHECK_HIP(hipMemcpy(collArg.outputCpu.ptr, collArg.outputGpu.ptr, numOutputBytes, hipMemcpyDeviceToHost)); @@ -816,9 +818,4 @@ namespace RcclUnitTesting if (this->verbose) INFO("Child %d finishes DestroyComms\n", this->childId); return TEST_SUCCESS; } - - ErrCode TestBedChild::Stop() - { - return TEST_SUCCESS; - } } diff --git a/projects/rccl/test/common/TestBedChild.hpp b/projects/rccl/test/common/TestBedChild.hpp index 0394fd170a..9e94c09b25 100644 --- a/projects/rccl/test/common/TestBedChild.hpp +++ b/projects/rccl/test/common/TestBedChild.hpp @@ -107,8 +107,5 @@ namespace RcclUnitTesting // Destroys RCCL communicators ErrCode DestroyComms(); - - // Stops this child process - ErrCode Stop(); }; } diff --git a/projects/rccl/test/common/main.cpp b/projects/rccl/test/common/main.cpp index b175982551..2c0597a61f 100644 --- a/projects/rccl/test/common/main.cpp +++ b/projects/rccl/test/common/main.cpp @@ -18,6 +18,8 @@ int main(int argc, char **argv) RcclUnitTesting::EnvVars ev; if (ev.showTiming) { + size_t totalTimeMsec = 0; + fflush(stdout); printf("[ TIMING ] %-20s: %-20s: %10s ms (%s)\n", "TEST SUITE", "TEST NAME", "TIME", "STATUS"); auto unitTest = ::testing::UnitTest::GetInstance(); for (int i = 0; i < unitTest->total_test_suite_count(); i++) @@ -31,12 +33,12 @@ int main(int argc, char **argv) if (!testInfo->should_run()) continue; auto testResult = testInfo->result(); if (testResult->Skipped()) continue; - - printf("[ TIMING ] %-20s: %-20s: %10ld ms (%4s)\n", testInfo->test_suite_name(), testInfo->name(), testResult->elapsed_time(), testResult->Passed() ? "PASS" : "FAIL"); + printf("[ TIMING ] %-20s: %-20s: %10.2f sec (%4s)\n", testInfo->test_suite_name(), testInfo->name(), testResult->elapsed_time() / 1000.0, testResult->Passed() ? "PASS" : "FAIL"); } - printf("[ TIMING ] %-20s: %-20s: %10ld ms (%4s)\n", suiteInfo->name(), "TOTAL", suiteInfo->elapsed_time(), suiteInfo->Passed() ? "PASS" : "FAIL"); + printf("[ TIMING ] %-20s: %-20s: %10.2f sec (%4s)\n", suiteInfo->name(), "TOTAL", suiteInfo->elapsed_time() / 1000.0, suiteInfo->Passed() ? "PASS" : "FAIL"); + totalTimeMsec += suiteInfo->elapsed_time(); } + printf("[ TIMING ] Total time: %10.2f minutes\n", totalTimeMsec / (60 * 1000.0)); } - return retCode; }