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

Graph unit tests (#656)

Browse Source 
* Adding hipGraph unit tests
This commit is contained in:
gilbertlee-amd
2022-12-01 10:28:42 -07:00
committed by GitHub
parent aebed537a5
commit faed69f9fc
31 changed files with 399 additions and 251 deletions
Download Patch File Download Diff File Expand all files Collapse all files
test/common/TestBedChild.cpp
+83 -16
View File
@@ -170,16 +170,16 @@ namespace RcclUnitTesting
if (useMultiRankPerGpu)
{
if (ncclCommInitRankMulti(&this->comms[localRank], this->totalRanks, id, globalRank, globalRank) != ncclSuccess)
if (ncclCommInitRankMulti(&this->comms[localRank], this->totalRanks, id, globalRank, globalRank) != ncclSuccess)
{
ERROR("Rank %d on child %d unable to call ncclCommInitRankMulti\n", globalRank, this->childId);
status = TEST_FAIL;
break;
}
ERROR("Rank %d on child %d unable to call ncclCommInitRankMulti\n", globalRank, this->childId);
status = TEST_FAIL;
break;
}
}
else
{
if (ncclCommInitRank(&this->comms[localRank], this->totalRanks, id, globalRank) != ncclSuccess)
if (ncclCommInitRank(&this->comms[localRank], this->totalRanks, id, globalRank) != ncclSuccess)
{
ERROR("Rank %d on child %d unable to call ncclCommInitRank\n", globalRank, this->childId);
status = TEST_FAIL;
@@ -337,6 +337,9 @@ namespace RcclUnitTesting
ErrCode TestBedChild::ExecuteCollectives()
{
bool useHipGraph = false;
PIPE_READ(useHipGraph);
int numRanksToExecute, tempRank;
std::vector<int> ranksToExecute = {};
PIPE_READ(numRanksToExecute);
@@ -345,7 +348,31 @@ namespace RcclUnitTesting
PIPE_READ(tempRank);
ranksToExecute.push_back(tempRank - this->rankOffset);
}
if (this->verbose) INFO("Child %d begins ExecuteCollectives()\n", this->childId);
if (this->verbose) INFO("Child %d begins ExecuteCollectives() %s\n", this->childId, useHipGraph ? "(using hipGraphs)" : "");
// Determine which local ranks to execute on
std::vector<int> localRanksToExecute;
for (int localRank = 0; localRank < this->deviceIds.size(); ++localRank)
{
// If ranksToExeute is empty, execute all local ranks belonging to this child
if (!ranksToExecute.empty() &&
(std::count(ranksToExecute.begin(), ranksToExecute.end(), localRank) == 0)) continue;
localRanksToExecute.push_back(localRank);
}
numRanksToExecute = (int)localRanksToExecute.size();
hipGraph_t graphs[numRanksToExecute];
hipGraphExec_t graphExec[numRanksToExecute];
// Start HIP graph stream capture if requested
if (useHipGraph)
{
for (int localRank : localRanksToExecute)
{
if (this->verbose) INFO("Capturing stream for rank %d\n", localRank);
CHECK_HIP(hipStreamBeginCapture(this->streams[localRank], hipStreamCaptureModeGlobal));
}
}
// Start group call
CHILD_NCCL_CALL(ncclGroupStart(), "ncclGroupStart");
@@ -354,16 +381,13 @@ namespace RcclUnitTesting
for (int collId = 0; collId < this->numCollectivesInGroup; ++collId)
{
// Loop over all local ranks
for (int localRank = 0; localRank < this->deviceIds.size(); ++localRank)
for (int localRank : localRanksToExecute)
{
// If ranks to execute is empty, execute all ranks belonging to child
if (!ranksToExecute.empty() && (std::count(ranksToExecute.begin(), ranksToExecute.end(), localRank) == 0)) continue;
CHECK_HIP(hipSetDevice(this->deviceIds[localRank]));
CollectiveArgs const& collArg = this->collArgs[localRank][collId];
if (this->printValues)
if (this->printValues && !useHipGraph)
{
int const numInputElementsToPrint = (this->printValues < 0 ? collArg.numInputElements : this->printValues);
PtrUnion inputCpu;
@@ -502,17 +526,60 @@ namespace RcclUnitTesting
// End group call
CHILD_NCCL_CALL(ncclGroupEnd(), "ncclGroupEnd");
// Synchronize
if (this->verbose) INFO("Child %d submits group call. Waiting for completion\n", this->childId);
for (int localRank = 0; localRank < this->streams.size(); ++localRank)
// Instantiate and launch HIP graph if requested
if (useHipGraph)
{
for (int localRank : localRanksToExecute)
{
if (this->verbose) INFO("Ending stream capture for rank %d\n", localRank);
CHECK_HIP(hipStreamEndCapture(this->streams[localRank], &graphs[localRank]));
if (this->verbose)
{
size_t numNodes;
hipGraphNode_t* nodes;
CHECK_HIP(hipGraphGetNodes(graphs[localRank], nodes, &numNodes));
INFO("Graph for rank %d has %lu nodes\n", localRank, numNodes);
}
if (this->verbose) INFO("Instantiating executable graph for rank %d\n", localRank);
CHECK_HIP(hipGraphInstantiate(&graphExec[localRank], graphs[localRank], NULL, NULL, 0));
}
for (int localRank : localRanksToExecute)
{
if (this->verbose) INFO("Launch graph for rank %d\n", localRank);
CHECK_HIP(hipGraphLaunch(graphExec[localRank], this->streams[localRank]));
}
}
else
{
if (this->verbose)
INFO("Child %d submits group call. Waiting for completion\n", this->childId);
}
// Synchronize
for (int localRank : localRanksToExecute)
{
if (this->verbose) INFO("Starting synchronization for rank %d\n", localRank);
CHECK_HIP(hipStreamSynchronize(this->streams[localRank]));
}
// Destroy graphs
if (useHipGraph)
{
for (int localRank : localRanksToExecute)
{
if (this->verbose) INFO("Destroying graphs for rank %d\n", localRank);
CHECK_HIP(hipGraphDestroy(graphs[localRank]));
CHECK_HIP(hipGraphExecDestroy(graphExec[localRank]));
}
}
if (this->printValues)
{
for (int collId = 0; collId < this->numCollectivesInGroup; ++collId)
for (int localRank = 0; localRank < this->deviceIds.size(); ++localRank)
for (int localRank : localRanksToExecute)
{
CollectiveArgs const& collArg = this->collArgs[localRank][collId];