AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Unit test performance refactor (#700)

Browse Source 
* Refactoring unit tests to improve performance
* Spawning child processes during InitComms instead of on TestBed construction
* Temporarily disabling graph unit tests
This commit is contained in:
gilbertlee-amd
2023-04-06 12:28:53 -06:00
committed by GitHub
parent 9fe5a349f1
commit 27e0cb43c2
52 changed files with 1597 additions and 1383 deletions
Download Patch File Download Diff File Expand all files Collapse all files
test/common/EnvVars.cpp
+23
View File
@@ -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"),
test/common/EnvVars.hpp
+7 -2
View File
@@ -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);
test/common/TestBed.cpp
+91 -47
View File
@@ -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)
{
test/common/TestBedChild.cpp
+4 -7
View File
@@ -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;
}
}
test/common/TestBedChild.hpp
-3
View File
@@ -107,8 +107,5 @@ namespace RcclUnitTesting
// Destroys RCCL communicators
ErrCode DestroyComms();
// Stops this child process
ErrCode Stop();
};
}
test/common/main.cpp
+6 -4
View File
@@ -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;
}