Unit test performance refactor (#700)
* Refactoring unit tests to improve performance
* Spawning child processes during InitComms instead of on TestBed construction
* Temporarily disabling graph unit tests
[ROCm/rccl commit: 27e0cb43c2]
Цей коміт міститься в:
gilbertlee-amd
GitHub
@@ -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)
|
||||
|
||||
@@ -37,7 +37,7 @@ def runCI =
|
||||
{
|
||||
platform, project->
|
||||
|
||||
commonGroovy.runTestCommand(platform, project, "*")
|
||||
commonGroovy.runTestCommand(platform, project, "-*Graph", "")
|
||||
}
|
||||
|
||||
def packageCommand =
|
||||
|
||||
@@ -37,7 +37,7 @@ def runCI =
|
||||
{
|
||||
platform, project->
|
||||
|
||||
commonGroovy.runTestCommand(platform, project, "-*ManagedMem")
|
||||
commonGroovy.runTestCommand(platform, project, "-*Graph", "UT_POW2_GPUS=1")
|
||||
}
|
||||
|
||||
def packageCommand =
|
||||
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 500};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllGather, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclBfloat16, ncclFloat64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {586};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllGather, InPlace)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {104857, 264};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllGather, InPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {958};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllGather, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1039203, 2500};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllGather, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {896};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {393216, 384};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllReduce, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat16, ncclFloat64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMin};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {12888};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllReduce, InPlace)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt32, ncclInt8};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclProd};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {384};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllReduce, InPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMax};
|
||||
std::vector<int> const roots = {0, 1};
|
||||
std::vector<int> const numElements = {393216};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllReduce, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {2500};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllReduce, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {4314};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> 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<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = testBed.GetAllSupportedDataTypes();
|
||||
std::vector<ncclRedOp_t> const redOps = testBed.GetAllSupportedRedOps();
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 1024};
|
||||
std::vector<bool> const inPlaceList = {false, true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> 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<ncclDataType_t> const& dataTypes = {ncclFloat32};
|
||||
ncclRedOp_t const redOp = ncclSum;
|
||||
std::vector<int> 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();
|
||||
}
|
||||
}
|
||||
@@ -1,33 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
#include "TestBed.hpp"
|
||||
#include <cstdlib>
|
||||
namespace RcclUnitTesting
|
||||
{
|
||||
TEST(AllReduce, DISABLED_Clique)
|
||||
{
|
||||
// Set clique env var prior to TestBed
|
||||
setenv("RCCL_ENABLE_CLIQUE", "1", 1);
|
||||
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = testBed.GetAllSupportedDataTypes();
|
||||
std::vector<ncclRedOp_t> const redOps = testBed.GetAllSupportedRedOps();
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false, true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
|
||||
unsetenv("RCCL_ENABLE_CLIQUE");
|
||||
}
|
||||
}
|
||||
@@ -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<ncclDataType_t> const& dataTypes = {ncclFloat};
|
||||
std::vector<ncclRedOp_t> const& redOps = {ncclSum};
|
||||
std::vector<int> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclProd};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclUint8, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclMax};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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 <cstdlib>
|
||||
|
||||
#include "TestBed.hpp"
|
||||
|
||||
extern "C" bool mscclUnitTestMode() { return true; }
|
||||
|
||||
namespace RcclUnitTesting
|
||||
{
|
||||
TEST(AllReduce, MscclSingleCall)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclProd};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {384 * 32 * 32, 384 * 32, 384};
|
||||
std::vector<bool> const inPlaceList = {true, false};
|
||||
std::vector<bool> const managedMemList = {true, false};
|
||||
std::vector<bool> 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<ncclDataType_t> const& dataTypes = {ncclFloat};
|
||||
std::vector<ncclRedOp_t> const& redOps = {ncclSum};
|
||||
std::vector<int> 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<ncclDataType_t> const& dataTypes = {ncclInt32, ncclFloat32, ncclFloat64};
|
||||
ncclRedOp_t const redOp = ncclSum;
|
||||
std::vector<int> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclDataType_t> const& dataTypes = {ncclFloat};
|
||||
std::vector<ncclRedOp_t> const& redOps = {ncclSum};
|
||||
std::vector<int> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat16, ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclMin};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclDataType_t> const& dataTypes = {ncclInt32, ncclFloat32, ncclFloat64};
|
||||
ncclRedOp_t const redOp = ncclSum;
|
||||
std::vector<int> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllToAll};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllToAll, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllToAll};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {5685};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllToAll, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllToAll};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {384 * 1024, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(AllToAll, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllToAll};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<size_t>& numInputElements,
|
||||
std::vector<size_t>& 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<ncclDataType_t> 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<size_t> numInputElements;
|
||||
std::vector<size_t> 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<ncclDataType_t> 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<size_t> numInputElements;
|
||||
std::vector<size_t> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllToAll};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllToAll};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclDataType_t> const& dataTypes = {ncclInt32, ncclFloat64, ncclFloat16};
|
||||
std::vector<int> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 500};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Broadcast, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclBfloat16, ncclFloat64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {586};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Broadcast, InPlace)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {104857, 264};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Broadcast, InPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {958};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Broadcast, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1039203, 2500};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Broadcast, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {896};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollBroadcast};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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()
|
||||
|
||||
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {1048576, 127};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Gather, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat16, ncclFloat64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {924873};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
|
||||
TEST(Gather, InPlace)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {13576};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Gather, InPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt64, ncclFloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Gather, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {1051234};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Gather, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {5231};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64, ncclFloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollAllReduce, ncclCollAllReduce, ncclCollAllReduce};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclSum, ncclSum};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat, ncclFloat, ncclFloat};
|
||||
std::vector<int> 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<ncclFunc_t> 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();
|
||||
}
|
||||
}
|
||||
@@ -1,61 +0,0 @@
|
||||
/*************************************************************************
|
||||
* Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* See LICENSE.txt for license information
|
||||
************************************************************************/
|
||||
#include "TestBed.hpp"
|
||||
#include <cstdlib>
|
||||
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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMax};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {393216, 384};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(ReduceScatter, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMax};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(ReduceScatter, InPlace)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclProd};
|
||||
std::vector<int> const roots = {0, 1};
|
||||
std::vector<int> const numElements = {542357};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(ReduceScatter, InPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclFloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMin};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {246};;
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(ReduceScatter, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt64, ncclUint8};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclAvg};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(ReduceScatter, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclAvg};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {6485423};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclProd};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};;
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclAvg};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMax};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {393216, 384};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Reduce, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat16, ncclFloat64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMin};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {393216};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Reduce, InPlace)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt32, ncclInt8};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclProd};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {384};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Reduce, InPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMax};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {393216};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Reduce, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {3524082, 2500};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Reduce, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {4314};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclProd};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMin, ncclMax, ncclAvg};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {393216, 384};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Scatter, OutOfPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {24658};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Scatter, InPlace)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt32};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Scatter, InPlaceGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclFloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {356};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Scatter, ManagedMem)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt64, ncclUint8};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {948576};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
|
||||
TEST(Scatter, ManagedMemGraph)
|
||||
{
|
||||
TestBed testBed;
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {125};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclUint8, ncclUint32, ncclUint64};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {true};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
@@ -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<ncclFunc_t> const funcTypes = {ncclCollScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {1};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
std::vector<bool> const useHipGraphList = {false, true};
|
||||
|
||||
testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
|
||||
inPlaceList, managedMemList, useHipGraphList);
|
||||
testBed.Finalize();
|
||||
}
|
||||
}
|
||||
+1
-1
@@ -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
|
||||
************************************************************************/
|
||||
@@ -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<ncclRedOp_t> const& EnvVars::GetAllSupportedRedOps()
|
||||
@@ -136,6 +148,16 @@ namespace RcclUnitTesting
|
||||
return dataTypes;
|
||||
}
|
||||
|
||||
std::vector<int> const& EnvVars::GetNumGpusList()
|
||||
{
|
||||
return numGpusList;
|
||||
}
|
||||
|
||||
std::vector<int> 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"),
|
||||
|
||||
@@ -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<ncclRedOp_t> const& GetAllSupportedRedOps();
|
||||
std::vector<ncclDataType_t> const& GetAllSupportedDataTypes();
|
||||
|
||||
std::vector<int> const& GetNumGpusList();
|
||||
std::vector<int> const& GetIsMultiProcessList();
|
||||
static void ShowConfig();
|
||||
|
||||
protected:
|
||||
std::vector<ncclRedOp_t> redOps; // Supported reduction ops [UT_REDOPS]
|
||||
std::vector<ncclDataType_t> dataTypes; // Support datatypes [UT_DATATYPES]
|
||||
std::vector<ncclRedOp_t> redOps; // Supported reduction ops [UT_REDOPS]
|
||||
std::vector<ncclDataType_t> dataTypes; // Support datatypes [UT_DATATYPES]
|
||||
std::vector<int> numGpusList; // List of # Gpus to use [UT_MIN_GPUS/UT_MAX_GPUS/UT_POW2_GPUS]
|
||||
std::vector<int> isMultiProcessList; // Single or multi process [UT_PROCESS_MASK]
|
||||
|
||||
// Helper functions to parse environment variables
|
||||
int GetEnvVar(std::string const varname, int defaultValue);
|
||||
|
||||
@@ -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<std::vector<int>> 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<int> 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<int> currentRanksEmpty = {};
|
||||
this->ExecuteCollectives(currentRanksEmpty, useHipGraphList[hgIdx]);
|
||||
if (testing::Test::HasFailure()) continue;
|
||||
if (testing::Test::HasFailure())
|
||||
{
|
||||
isCorrect = false;
|
||||
continue;
|
||||
}
|
||||
this->ValidateResults(isCorrect);
|
||||
if (!isCorrect)
|
||||
{
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -107,8 +107,5 @@ namespace RcclUnitTesting
|
||||
|
||||
// Destroys RCCL communicators
|
||||
ErrCode DestroyComms();
|
||||
|
||||
// Stops this child process
|
||||
ErrCode Stop();
|
||||
};
|
||||
}
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
|
||||
Посилання в новій задачі
Заблокувати користувача