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

SWDEV-240806 - hipGraph performance create new graph commands for every launch

Browse Source 
Change-Id: I7360b48055f67a91c23169ca532a19124d40e98b


[ROCm/clr commit: ec6b8d82b5]
This commit is contained in:
anusha GodavarthySurya
2021-11-17 00:13:52 -08:00
parent 417851d666
commit abb6f92138
3 changed files with 564 additions and 282 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_graph.cpp
+53 -32
View File
@@ -25,7 +25,6 @@
#include "hip_event.hpp"
thread_local std::vector<hipStream_t> g_captureStreams;
std::unordered_map<amd::Command*, hipGraphExec_t> hipGraphExec::activeGraphExec_;
inline void ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph,
const hipGraphNode_t* pDependencies, size_t numDependencies) {
@@ -419,16 +418,18 @@ hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
hipError_t hipGraphMemcpyNodeSetParams1D(hipGraphNode_t node, void* dst, const void* src,
size_t count, hipMemcpyKind kind) {
HIP_INIT_API(hipGraphMemcpyNodeSetParams1D, node, dst, src, count, kind);
reinterpret_cast<hipGraphMemcpyNode1D*>(node)->SetParams(dst, src, count, kind);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNode1D*>(node)->SetParams(dst, src, count, kind));
}
hipError_t hipGraphExecMemcpyNodeSetParams1D(hipGraphExec_t hGraphExec, hipGraphNode_t node,
void* dst, const void* src, size_t count,
hipMemcpyKind kind) {
HIP_INIT_API(hipGraphExecMemcpyNodeSetParams1D, hGraphExec, node, dst, src, count, kind);
HIP_RETURN(
reinterpret_cast<hipGraphMemcpyNode1D*>(node)->SetCommandParams(dst, src, count, kind));
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node);
if (clonedNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNode1D*>(clonedNode)->SetParams(dst, src, count, kind));
}
hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
@@ -468,12 +469,14 @@ hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t grap
hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize) {
std::unordered_map<Node, Node> clonedNodes;
hipGraph_t clonedGraph = graph->clone(clonedNodes);
std::vector<std::vector<Node>> parallelLists;
std::unordered_map<Node, std::vector<Node>> nodeWaitLists;
graph->GetRunList(parallelLists, nodeWaitLists);
clonedGraph->GetRunList(parallelLists, nodeWaitLists);
std::vector<Node> levelOrder;
graph->LevelOrder(levelOrder);
*pGraphExec = new hipGraphExec(levelOrder, parallelLists, nodeWaitLists);
clonedGraph->LevelOrder(levelOrder);
*pGraphExec = new hipGraphExec(levelOrder, parallelLists, nodeWaitLists, clonedNodes);
if (*pGraphExec != nullptr) {
return (*pGraphExec)->Init();
} else {
@@ -502,7 +505,7 @@ hipError_t ihipGraphlaunch(hipGraphExec_t graphExec, hipStream_t stream) {
hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) {
HIP_INIT_API(hipGraphLaunch, graphExec, stream);
if (graphExec == nullptr || stream == nullptr) {
if (graphExec == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(ihipGraphlaunch(graphExec, stream));
@@ -569,8 +572,7 @@ hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node,
if (node == nullptr || pNodeParams == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
reinterpret_cast<hipGraphKernelNode*>(node)->SetParams(pNodeParams);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphKernelNode*>(node)->SetParams(pNodeParams));
}
hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams) {
@@ -587,14 +589,17 @@ hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DPar
if (node == nullptr || pNodeParams == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
reinterpret_cast<hipGraphMemcpyNode*>(node)->SetParams(pNodeParams);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNode*>(node)->SetParams(pNodeParams));
}
hipError_t hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
hipMemcpy3DParms* pNodeParams) {
HIP_INIT_API(hipGraphExecMemcpyNodeSetParams, hGraphExec, node, pNodeParams);
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNode*>(node)->SetCommandParams(pNodeParams));
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node);
if (clonedNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNode*>(clonedNode)->SetParams(pNodeParams));
}
hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams) {
@@ -611,8 +616,7 @@ hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParam
if (node == nullptr || pNodeParams == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
reinterpret_cast<hipGraphMemsetNode*>(node)->SetParams(pNodeParams);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphMemsetNode*>(node)->SetParams(pNodeParams));
}
hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
@@ -638,7 +642,11 @@ hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo
if (hGraphExec == nullptr || node == nullptr || pNodeParams == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphKernelNode*>(node)->SetCommandParams(pNodeParams));
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node);
if (clonedNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphKernelNode*>(clonedNode)->SetParams(pNodeParams));
}
hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph) {
@@ -879,9 +887,8 @@ hipError_t hipGraphAddMemcpyNodeFromSymbol(hipGraphNode_t* pGraphNode, hipGraph_
hipError_t hipGraphMemcpyNodeSetParamsFromSymbol(hipGraphNode_t node, void* dst, const void* symbol,
size_t count, size_t offset, hipMemcpyKind kind) {
HIP_INIT_API(hipGraphMemcpyNodeSetParamsFromSymbol, node, dst, symbol, count, offset, kind);
reinterpret_cast<hipGraphMemcpyNodeFromSymbol*>(node)->SetParams(dst, symbol, count, offset,
kind);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeFromSymbol*>(node)->SetParams(dst, symbol, count,
offset, kind));
}
hipError_t hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node,
@@ -889,8 +896,12 @@ hipError_t hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec,
size_t offset, hipMemcpyKind kind) {
HIP_INIT_API(hipGraphExecMemcpyNodeSetParamsFromSymbol, hGraphExec, node, dst, symbol, count,
offset, kind);
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeFromSymbol*>(node)->SetCommandParams(
dst, symbol, count, offset, kind));
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node);
if (clonedNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeFromSymbol*>(clonedNode)
->SetParams(dst, symbol, count, offset, kind));
}
hipError_t hipGraphAddMemcpyNodeToSymbol(hipGraphNode_t* pGraphNode, hipGraph_t graph,
@@ -915,8 +926,8 @@ hipError_t hipGraphMemcpyNodeSetParamsToSymbol(hipGraphNode_t node, const void*
const void* src, size_t count, size_t offset,
hipMemcpyKind kind) {
HIP_INIT_API(hipGraphMemcpyNodeSetParamsToSymbol, symbol, src, count, offset, kind);
reinterpret_cast<hipGraphMemcpyNodeToSymbol*>(node)->SetParams(symbol, src, count, offset, kind);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeToSymbol*>(node)->SetParams(symbol, src, count,
offset, kind));
}
@@ -926,8 +937,12 @@ hipError_t hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hi
hipMemcpyKind kind) {
HIP_INIT_API(hipGraphExecMemcpyNodeSetParamsToSymbol, hGraphExec, node, symbol, src, count,
offset, kind);
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeToSymbol*>(node)->SetCommandParams(
symbol, src, count, offset, kind));
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node);
if (clonedNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeToSymbol*>(clonedNode)
->SetParams(symbol, src, count, offset, kind));
}
hipError_t hipGraphAddEventRecordNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
@@ -957,8 +972,7 @@ hipError_t hipGraphEventRecordNodeSetEvent(hipGraphNode_t node, hipEvent_t event
if (node == nullptr || event == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
reinterpret_cast<hipGraphEventRecordNode*>(node)->SetParams(event);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphEventRecordNode*>(node)->SetParams(event));
}
hipError_t hipGraphExecEventRecordNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
@@ -967,7 +981,11 @@ hipError_t hipGraphExecEventRecordNodeSetEvent(hipGraphExec_t hGraphExec, hipGra
if (hGraphExec == nullptr || hNode == nullptr || event == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphEventRecordNode*>(hNode)->SetExecParams(event));
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode);
if (clonedNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphEventRecordNode*>(clonedNode)->SetParams(event));
}
hipError_t hipGraphAddEventWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
@@ -996,8 +1014,7 @@ hipError_t hipGraphEventWaitNodeSetEvent(hipGraphNode_t node, hipEvent_t event)
if (node == nullptr || event == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
reinterpret_cast<hipGraphEventWaitNode*>(node)->SetParams(event);
HIP_RETURN(hipSuccess);
HIP_RETURN(reinterpret_cast<hipGraphEventWaitNode*>(node)->SetParams(event));
}
hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
@@ -1006,5 +1023,9 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph
if (hGraphExec == nullptr || hNode == nullptr || event == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphEventRecordNode*>(hNode)->SetExecParams(event));
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode);
if (clonedNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphEventRecordNode*>(clonedNode)->SetParams(event));
}
projects/clr/hipamd/src/hip_graph_internal.cpp
+275 -110
View File
@@ -46,6 +46,40 @@ const char* GetGraphNodeTypeString(uint32_t op) {
};
int hipGraphNode::nextID = 0;
hipError_t hipGraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size_t count,
hipMemcpyKind kind) {
hipError_t status = ihipMemcpy_validate(dst, src, count, kind);
if (status != hipSuccess) {
return status;
}
size_t sOffsetOrig = 0;
amd::Memory* origSrcMemory = getMemoryObject(src, sOffsetOrig);
size_t dOffsetOrig = 0;
amd::Memory* origDstMemory = getMemoryObject(dst, dOffsetOrig);
size_t sOffset = 0;
amd::Memory* srcMemory = getMemoryObject(src, sOffset);
size_t dOffset = 0;
amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
if (origDstMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
return hipErrorInvalidValue;
}
} else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
if (origSrcMemory->getContext().devices()[0] != srcMemory->getContext().devices()[0]) {
return hipErrorInvalidValue;
}
} else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
if (origDstMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
return hipErrorInvalidValue;
}
if (origSrcMemory->getContext().devices()[0] != srcMemory->getContext().devices()[0]) {
return hipErrorInvalidValue;
}
}
return hipSuccess;
}
hipError_t hipGraphMemcpyNode1D::SetCommandParams(void* dst, const void* src, size_t count,
hipMemcpyKind kind) {
@@ -54,9 +88,9 @@ hipError_t hipGraphMemcpyNode1D::SetCommandParams(void* dst, const void* src, si
return status;
}
size_t sOffsetOrig = 0;
amd::Memory* origSrcMemory = getMemoryObject(src_, sOffsetOrig);
amd::Memory* origSrcMemory = getMemoryObject(src, sOffsetOrig);
size_t dOffsetOrig = 0;
amd::Memory* origDstMemory = getMemoryObject(dst_, dOffsetOrig);
amd::Memory* origDstMemory = getMemoryObject(dst, dOffsetOrig);
size_t sOffset = 0;
amd::Memory* srcMemory = getMemoryObject(src, sOffset);
@@ -97,6 +131,155 @@ hipError_t hipGraphMemcpyNode1D::SetCommandParams(void* dst, const void* src, si
return hipSuccess;
}
hipError_t hipGraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParams) {
hipError_t status = ihipMemcpy3D_validate(pNodeParams);
if (status != hipSuccess) {
return status;
}
const HIP_MEMCPY3D pCopy = hip::getDrvMemcpy3DDesc(*pNodeParams);
// If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the
// (unified virtual address space) base address of the source data and the bytes per row to apply.
// {src/dst}Array is ignored.
hipMemoryType srcMemoryType = pCopy.srcMemoryType;
if (srcMemoryType == hipMemoryTypeUnified) {
srcMemoryType =
amd::MemObjMap::FindMemObj(pCopy.srcDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost;
if (srcMemoryType == hipMemoryTypeHost) {
// {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system
// memory.
const_cast<HIP_MEMCPY3D*>(&pCopy)->srcHost = pCopy.srcDevice;
}
}
hipMemoryType dstMemoryType = pCopy.dstMemoryType;
if (dstMemoryType == hipMemoryTypeUnified) {
dstMemoryType =
amd::MemObjMap::FindMemObj(pCopy.dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost;
if (srcMemoryType == hipMemoryTypeHost) {
const_cast<HIP_MEMCPY3D*>(&pCopy)->dstHost = pCopy.dstDevice;
}
}
// If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned.
// In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning.
if (srcMemoryType == hipMemoryTypeHost) {
amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.srcHost);
srcMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost;
if (srcMemoryType == hipMemoryTypeDevice) {
const_cast<HIP_MEMCPY3D*>(&pCopy)->srcDevice = const_cast<void*>(pCopy.srcHost);
}
}
if (dstMemoryType == hipMemoryTypeHost) {
amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.dstHost);
dstMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost;
if (dstMemoryType == hipMemoryTypeDevice) {
const_cast<HIP_MEMCPY3D*>(&pCopy)->dstDevice = const_cast<void*>(pCopy.dstDevice);
}
}
amd::Coord3D srcOrigin = {pCopy.srcXInBytes, pCopy.srcY, pCopy.srcZ};
amd::Coord3D dstOrigin = {pCopy.dstXInBytes, pCopy.dstY, pCopy.dstZ};
amd::Coord3D copyRegion = {pCopy.WidthInBytes, pCopy.Height, pCopy.Depth};
if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeDevice)) {
// Host to Device.
amd::Memory* dstMemory;
amd::BufferRect srcRect;
amd::BufferRect dstRect;
status =
ihipMemcpyHtoDValidate(pCopy.srcHost, pCopy.dstDevice, srcOrigin, dstOrigin, copyRegion,
pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, pCopy.dstPitch,
pCopy.dstPitch * pCopy.dstHeight, dstMemory, srcRect, dstRect);
if (status != hipSuccess) {
return status;
}
} else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeHost)) {
// Device to Host.
amd::Memory* srcMemory;
amd::BufferRect srcRect;
amd::BufferRect dstRect;
status =
ihipMemcpyDtoHValidate(pCopy.srcDevice, pCopy.dstHost, srcOrigin, dstOrigin, copyRegion,
pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, pCopy.dstPitch,
pCopy.dstPitch * pCopy.dstHeight, srcMemory, srcRect, dstRect);
if (status != hipSuccess) {
return status;
}
} else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeDevice)) {
// Device to Device.
amd::Memory* srcMemory;
amd::Memory* dstMemory;
amd::BufferRect srcRect;
amd::BufferRect dstRect;
status = ihipMemcpyDtoDValidate(pCopy.srcDevice, pCopy.dstDevice, srcOrigin, dstOrigin,
copyRegion, pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight,
pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcMemory,
dstMemory, srcRect, dstRect);
if (status != hipSuccess) {
return status;
}
} else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) {
amd::Image* dstImage;
amd::BufferRect srcRect;
status =
ihipMemcpyHtoAValidate(pCopy.srcHost, pCopy.dstArray, srcOrigin, dstOrigin, copyRegion,
pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, dstImage, srcRect);
if (status != hipSuccess) {
return status;
}
} else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) {
// Image to Host.
amd::Image* srcImage;
amd::BufferRect dstRect;
status =
ihipMemcpyAtoHValidate(pCopy.srcArray, pCopy.dstHost, srcOrigin, dstOrigin, copyRegion,
pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcImage, dstRect);
if (status != hipSuccess) {
return status;
}
} else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeArray)) {
// Device to Image.
amd::Image* dstImage;
amd::Memory* srcMemory;
amd::BufferRect dstRect;
amd::BufferRect srcRect;
status = ihipMemcpyDtoAValidate(pCopy.srcDevice, pCopy.dstArray, srcOrigin, dstOrigin,
copyRegion, pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight,
dstImage, srcMemory, dstRect, srcRect);
if (status != hipSuccess) {
return status;
}
} else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeDevice)) {
// Image to Device.
amd::BufferRect srcRect;
amd::BufferRect dstRect;
amd::Memory* dstMemory;
amd::Image* srcImage;
status = ihipMemcpyAtoDValidate(pCopy.srcArray, pCopy.dstDevice, srcOrigin, dstOrigin,
copyRegion, pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight,
dstMemory, srcImage, srcRect, dstRect);
if (status != hipSuccess) {
return status;
}
} else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeArray)) {
amd::Image* srcImage;
amd::Image* dstImage;
status = ihipMemcpyAtoAValidate(pCopy.srcArray, pCopy.dstArray, srcOrigin, dstOrigin,
copyRegion, srcImage, dstImage);
if (status != hipSuccess) {
return status;
}
} else {
ShouldNotReachHere();
}
return hipSuccess;
}
hipError_t hipGraphMemcpyNode::SetCommandParams(const hipMemcpy3DParms* pNodeParams) {
hipError_t status = ihipMemcpy3D_validate(pNodeParams);
if (status != hipSuccess) {
@@ -425,9 +608,8 @@ void ihipGraph::LevelOrder(std::vector<Node>& levelOrder) {
}
}
ihipGraph* ihipGraph::clone() const {
ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const {
ihipGraph* newGraph = new ihipGraph();
std::unordered_map<Node, Node> clonedNodes;
for (auto entry : vertices_) {
hipGraphNode* node = entry->clone();
newGraph->vertices_.push_back(node);
@@ -455,9 +637,14 @@ ihipGraph* ihipGraph::clone() const {
return newGraph;
}
hipError_t hipGraphExec::CreateQueues() {
parallelQueues_.reserve(parallelLists_.size());
for (size_t i = 0; i < parallelLists_.size(); i++) {
ihipGraph* ihipGraph::clone() const {
std::unordered_map<Node, Node> clonedNodes;
return clone(clonedNodes);
}
hipError_t hipGraphExec::CreateQueues(size_t numQueues) {
parallelQueues_.reserve(numQueues);
for (size_t i = 0; i < numQueues; i++) {
amd::HostQueue* queue;
cl_command_queue_properties properties =
(callbacks_table.is_enabled() || HIP_FORCE_QUEUE_PROFILING) ? CL_QUEUE_PROFILING_ENABLE : 0;
@@ -476,80 +663,44 @@ hipError_t hipGraphExec::CreateQueues() {
return hipSuccess;
}
hipError_t hipGraphExec::FillCommands() {
// Create commands
int i = 0;
hipError_t hipGraphExec::Init() {
hipError_t status;
for (const auto& list : parallelLists_) {
for (auto& node : list) {
status = node->CreateCommand(parallelQueues_[i]);
if (status != hipSuccess) return status;
}
i++;
size_t reqNumQueues = 0;
for (auto& node : levelOrder_) {
reqNumQueues += node->GetNumParallelQueues();
}
status = CreateQueues(parallelLists_.size() - 1 + reqNumQueues);
return status;
}
i = 0;
// For nodes that has embedded child graph
for (const auto& list : parallelLists_) {
for (auto& node : list) {
node->UpdateEventWaitLists();
}
i++;
}
// Add waitlists for all the commands
for (auto entry : nodeWaitLists_) {
hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
std::unordered_map<Node, std::vector<Node>>& nodeWaitLists,
std::vector<Node>& levelOrder, amd::Command*& rootCommand,
amd::Command*& endCommand, amd::HostQueue* queue) {
hipError_t status;
for (auto& node : levelOrder) {
// TODO: clone commands from next launch
status = node->CreateCommand(node->GetQueue());
if (status != hipSuccess) return status;
amd::Command::EventWaitList waitList;
for (auto depNode : entry.second) {
for (auto depNode : nodeWaitLists[node]) {
for (auto command : depNode->GetCommands()) {
waitList.push_back(command);
}
}
for (auto command : entry.first->GetCommands()) {
command->updateEventWaitList(waitList);
node->UpdateEventWaitLists(waitList);
}
// rootCommand ensures graph is started (all parallel branches) after all the previous work is
// finished
bool first = true;
for (auto& singleList : parallelLists) {
if (first) {
first = false;
continue;
}
}
return status;
}
hipError_t hipGraphExec::Init() {
hipError_t status;
status = CreateQueues();
if (status != hipSuccess) {
return status;
}
status = FillCommands();
if (status != hipSuccess) {
return status;
}
rootCommand_ = nullptr;
/// stream should execute next command after graph finishes
/// Add marker to the stream that waits for all the last commands in parallel queues of graph
for (auto& singleList : parallelLists_) {
graphLastCmdWaitList_.push_back(singleList.back()->GetCommands().back());
}
return status;
}
void hipGraphExec::ResetGraph(cl_event event, cl_int command_exec_status, void* user_data) {
ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Inside resetGraph!\n");
hipGraphExec_t graphExec =
hipGraphExec::activeGraphExec_[reinterpret_cast<amd::Command*>(user_data)];
if (graphExec != nullptr) {
for (auto& node : graphExec->levelOrder_) {
node->ResetStatus();
}
graphExec->rootCommand_->resetStatus(CL_INT_MAX);
graphExec->bExecPending_.store(false);
} else {
ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] graphExec is nullptr during resetGraph!\n");
}
}
void hipGraphExec::UpdateGraphToWaitOnRoot() {
for (auto& singleList : parallelLists_) {
rootCommand = new amd::Marker(*queue, false, {});
amd::Command::EventWaitList waitList;
waitList.push_back(rootCommand_);
waitList.push_back(rootCommand);
if (!singleList.empty()) {
auto commands = singleList[0]->GetCommands();
if (!commands.empty()) {
@@ -557,56 +708,70 @@ void hipGraphExec::UpdateGraphToWaitOnRoot() {
}
}
}
// endCommand ensures next enqueued ones start after graph is finished (all parallel branches)
amd::Command::EventWaitList graphLastCmdWaitList;
first = true;
for (auto& singleList : parallelLists) {
if (first) {
first = false;
continue;
}
graphLastCmdWaitList.push_back(singleList.back()->GetCommands().back());
}
if (!graphLastCmdWaitList.empty()) {
endCommand = new amd::Marker(*queue, false, graphLastCmdWaitList);
if (endCommand == nullptr) {
return hipErrorOutOfMemory;
}
}
return hipSuccess;
}
void UpdateQueue(std::vector<std::vector<Node>>& parallelLists, amd::HostQueue*& queue,
hipGraphExec* ptr) {
int i = 0;
for (const auto& list : parallelLists) {
// first parallel list will be launched on the same queue as parent
if (i == 0) {
for (auto& node : list) {
node->SetQueue(queue, ptr);
}
} else { // New queue for parallel branches
amd::HostQueue* paralleQueue = ptr->GetAvailableQueue();
for (auto& node : list) {
node->SetQueue(paralleQueue, ptr);
}
}
i++;
}
}
hipError_t hipGraphExec::Run(hipStream_t stream) {
if (bExecPending_.load() == true) {
ClPrint(
amd::LOG_INFO, amd::LOG_CODE,
"[hipGraph] Same graph launched while previous one is active, wait for it to finish!\n");
lastEnqueuedGraphCmd_->awaitCompletion();
}
hipError_t status;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
return hipErrorInvalidResourceHandle;
}
if (rootCommand_ == nullptr || rootCommand_->queue() != queue) {
if (rootCommand_ != nullptr) {
rootCommand_->release();
}
rootCommand_ = new amd::Marker(*queue, false, {});
UpdateGraphToWaitOnRoot();
UpdateQueue(parallelLists_, queue, this);
amd::Command* rootCommand = nullptr;
amd::Command* endCommand = nullptr;
status =
FillCommands(parallelLists_, nodeWaitLists_, levelOrder_, rootCommand, endCommand, queue);
if (status != hipSuccess) {
return status;
}
if (rootCommand != nullptr) {
rootCommand->enqueue();
rootCommand->release();
}
rootCommand_->enqueue();
for (auto& node : levelOrder_) {
node->EnqueueCommands(stream);
}
amd::Command* command = new amd::Marker(*queue, false, graphLastCmdWaitList_);
if (command == nullptr) {
return hipErrorOutOfMemory;
if (endCommand != nullptr) {
endCommand->enqueue();
endCommand->release();
}
amd::Event& event = command->event();
if (!event.setCallback(CL_COMPLETE, hipGraphExec::ResetGraph, command)) {
return hipErrorInvalidHandle;
}
hipGraphExec::activeGraphExec_[command] = this;
lastEnqueuedGraphCmd_ = command;
bExecPending_.store(true);
command->enqueue();
// Add the new barrier to stall the stream, until the callback is done
amd::Command::EventWaitList eventWaitList;
eventWaitList.push_back(command);
amd::Command* block_command = new amd::Marker(*queue, !kMarkerDisableFlush, eventWaitList);
if (block_command == nullptr) {
return hipErrorInvalidValue;
}
block_command->enqueue();
block_command->release();
command->release();
ResetQueueIndex();
return hipSuccess;
}
projects/clr/hipamd/src/hip_graph_internal.hpp
+236 -140
View File
@@ -34,8 +34,15 @@
typedef hipGraphNode* Node;
hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams);
hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
std::unordered_map<Node, std::vector<Node>>& nodeWaitLists,
std::vector<Node>& levelOrder, amd::Command*& rootCommand,
amd::Command*& endCommand, amd::HostQueue* queue);
void UpdateQueue(std::vector<std::vector<Node>>& parallelLists, amd::HostQueue*& queue,
hipGraphExec* ptr);
struct hipGraphNode {
protected:
amd::HostQueue* queue_;
uint32_t level_;
unsigned int id_;
hipGraphNodeType type_;
@@ -69,6 +76,7 @@ struct hipGraphNode {
id_ = node.id_;
parentGraph_ = nullptr;
}
virtual ~hipGraphNode() {
for (auto node : edges_) {
node->RemoveDependency(this);
@@ -77,8 +85,16 @@ struct hipGraphNode {
node->RemoveEdge(this);
}
}
amd::HostQueue* GetQueue() { return queue_; }
virtual void SetQueue(amd::HostQueue* queue, hipGraphExec* ptr = nullptr) { queue_ = queue; }
/// Create amd::command for the graph node
virtual hipError_t CreateCommand(amd::HostQueue* queue) { return hipSuccess; }
virtual hipError_t CreateCommand(amd::HostQueue* queue) {
commands_.clear();
queue_ = queue;
return hipSuccess;
}
/// Method to release amd::command part of node
virtual void ReleaseCommand() {
for (auto command : commands_) {
@@ -89,7 +105,7 @@ struct hipGraphNode {
/// Return node unique ID
int GetID() const { return id_; }
/// Returns command for graph node
std::vector<amd::Command*>& GetCommands() { return commands_; }
virtual std::vector<amd::Command*>& GetCommands() { return commands_; }
/// Returns graph node type
hipGraphNodeType GetType() const { return type_; }
/// Returns graph node in coming edges
@@ -158,17 +174,17 @@ struct hipGraphNode {
/// Get levelorder of the nodes embedded as part of the graphnode(e.g. ChildGraph)
virtual void LevelOrder(std::vector<Node>& levelOrder) {}
/// Update waitlist of the nodes embedded as part of the graphnode(e.g. ChildGraph)
virtual void UpdateEventWaitLists() {}
virtual void UpdateEventWaitLists(amd::Command::EventWaitList waitList) {
for (auto command : commands_) {
command->updateEventWaitList(waitList);
}
}
virtual size_t GetNumParallelQueues() { return 0; }
/// Enqueue commands part of the node
virtual void EnqueueCommands(hipStream_t stream) {
for (auto& command : commands_) {
command->enqueue();
}
}
/// Reset commands part of the node
virtual void ResetStatus() {
for (auto& command : commands_) {
command->resetStatus(CL_INT_MAX);
command->release();
}
}
ihipGraph* GetParentGraph() { return parentGraph_; }
@@ -206,17 +222,67 @@ struct ihipGraph {
void GetRunList(std::vector<std::vector<Node>>& parallelList,
std::unordered_map<Node, std::vector<Node>>& dependencies);
void LevelOrder(std::vector<Node>& levelOrder);
ihipGraph* clone(std::unordered_map<Node, Node>& clonedNodes) const;
ihipGraph* clone() const;
};
struct hipGraphExec {
std::vector<std::vector<Node>> parallelLists_;
// level order of the graph doesn't include nodes embedded as part of the child graph
std::vector<Node> levelOrder_;
std::unordered_map<Node, std::vector<Node>> nodeWaitLists_;
std::vector<amd::HostQueue*> parallelQueues_;
uint currentQueueIndex_;
std::unordered_map<Node, Node> clonedNodes_;
amd::Command* lastEnqueuedCommand_;
public:
hipGraphExec(std::vector<Node>& levelOrder, std::vector<std::vector<Node>>& lists,
std::unordered_map<Node, std::vector<Node>>& nodeWaitLists,
std::unordered_map<Node, Node>& clonedNodes)
: parallelLists_(lists),
levelOrder_(levelOrder),
nodeWaitLists_(nodeWaitLists),
clonedNodes_(clonedNodes),
lastEnqueuedCommand_(nullptr),
currentQueueIndex_(0) {}
~hipGraphExec() {
// new commands are launched for every launch they are destroyed as and when command is
// terminated after it complete execution
for (auto queue : parallelQueues_) {
queue->release();
}
for (auto it = clonedNodes_.begin(); it != clonedNodes_.end(); it++) delete it->second;
}
Node GetClonedNode(Node node) {
Node clonedNode;
if (clonedNodes_.find(node) == clonedNodes_.end()) {
return nullptr;
} else {
clonedNode = clonedNodes_[node];
}
return clonedNode;
}
amd::HostQueue* GetAvailableQueue() { return parallelQueues_[currentQueueIndex_++]; }
void ResetQueueIndex() { currentQueueIndex_ = 0; }
hipError_t Init();
hipError_t CreateQueues(size_t numQueues);
hipError_t Run(hipStream_t stream);
};
struct hipChildGraphNode : public hipGraphNode {
struct ihipGraph* childGraph_;
std::vector<Node> childGraphlevelOrder_;
std::vector<std::vector<Node>> parallelLists_;
std::unordered_map<Node, std::vector<Node>> nodeWaitLists_;
amd::Command* lastEnqueuedCommand_;
public:
hipChildGraphNode(ihipGraph* g) : hipGraphNode(hipGraphNodeTypeGraph) {
// ToDo: clone the child graph
childGraph_ = g->clone();
lastEnqueuedCommand_ = nullptr;
}
~hipChildGraphNode() { delete childGraph_; }
@@ -231,73 +297,70 @@ struct hipChildGraphNode : public hipGraphNode {
ihipGraph* GetChildGraph() { return childGraph_; }
hipError_t CreateCommand(amd::HostQueue* queue) {
commands_.reserve(2);
amd::Command::EventWaitList eventWaitList;
// Command for start of the graph
commands_.push_back(new amd::Marker(*queue, false, eventWaitList));
// Command for end of the graph
commands_.push_back(new amd::Marker(*queue, false, eventWaitList));
childGraph_->LevelOrder(childGraphlevelOrder_);
return hipSuccess;
}
void UpdateEventWaitLists() {
// ChildGraph should start after all parents
if (commands_.size() == 2) {
std::vector<Node> rootNodes = childGraph_->GetRootNodes();
amd::Command::EventWaitList waitList;
waitList.push_back(commands_[0]);
for (auto& node : rootNodes) {
for (auto command : node->GetCommands()) {
command->updateEventWaitList(waitList);
}
}
waitList.clear();
// End command should wait for graph to finish
std::vector<Node> leafNodes = childGraph_->GetLeafNodes();
for (auto& node : leafNodes) {
for (auto command : node->GetCommands()) {
waitList.push_back(command);
}
}
commands_[1]->updateEventWaitList(waitList);
} else {
ClPrint(amd::LOG_ERROR, amd::LOG_CODE,
"[hipGraph] childgraph node commands are not created!\n");
}
}
void ResetStatus() {
if (commands_.size() == 2) {
commands_[0]->resetStatus(CL_INT_MAX);
commands_[1]->resetStatus(CL_INT_MAX);
}
size_t GetNumParallelQueues() {
LevelOrder(childGraphlevelOrder_);
size_t num = 0;
for (auto& node : childGraphlevelOrder_) {
node->ResetStatus();
num += node->GetNumParallelQueues();
}
// returns total number of parallel queues required for child graph nodes to be launched
// first parallel list will be launched on the same queue as parent
return num + (parallelLists_.size() - 1);
}
void SetQueue(amd::HostQueue* queue, hipGraphExec* ptr = nullptr) {
queue_ = queue;
UpdateQueue(parallelLists_, queue, ptr);
}
// For nodes that are dependent on the child graph node waitlist is the last node of the first
// parallel list
std::vector<amd::Command*>& GetCommands() { return parallelLists_[0].back()->GetCommands(); }
// Create child graph node commands and set waitlists
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
commands_.reserve(2);
amd::Command* rootCommand = nullptr;
amd::Command* endCommand = nullptr;
status = FillCommands(parallelLists_, nodeWaitLists_, childGraphlevelOrder_, rootCommand,
endCommand, queue);
if (rootCommand != nullptr) {
commands_.push_back(rootCommand);
}
if (endCommand != nullptr) {
commands_.push_back(endCommand);
}
return status;
}
//
void UpdateEventWaitLists(amd::Command::EventWaitList waitList) {
parallelLists_[0].front()->UpdateEventWaitLists(waitList);
}
void GetRunList(std::vector<std::vector<Node>>& parallelList,
std::unordered_map<Node, std::vector<Node>>& dependencies) {
childGraph_->GetRunList(parallelList, dependencies);
childGraph_->GetRunList(parallelLists_, nodeWaitLists_);
}
void LevelOrder(std::vector<Node>& levelOrder) { childGraph_->LevelOrder(levelOrder); }
void EnqueueCommands(hipStream_t stream) {
if (commands_.size() == 2) {
// enqueue child graph start command
// enqueue child graph start command
if (commands_.size() == 1) {
commands_[0]->enqueue();
// enqueue nodes in child graph in level order
for (auto& node : childGraphlevelOrder_) {
node->EnqueueCommands(stream);
}
// enqueue child graph end command
}
// enqueue nodes in child graph in level order
for (auto& node : childGraphlevelOrder_) {
node->EnqueueCommands(stream);
}
// enqueue child graph end command
if (commands_.size() == 2) {
commands_[1]->enqueue();
} else {
ClPrint(amd::LOG_ERROR, amd::LOG_CODE,
"[hipGraph] childgraph node commands are not created!\n");
}
}
};
@@ -321,9 +384,13 @@ class hipGraphKernelNode : public hipGraphNode {
return new hipGraphKernelNode(static_cast<hipGraphKernelNode const&>(*this));
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
commands_.reserve(1);
amd::Command* command;
hipError_t status = ihipLaunchKernelCommand(
status = ihipLaunchKernelCommand(
command, func_, pKernelParams_->gridDim.x * pKernelParams_->blockDim.x,
pKernelParams_->gridDim.y * pKernelParams_->blockDim.y,
pKernelParams_->gridDim.z * pKernelParams_->blockDim.z, pKernelParams_->blockDim.x,
@@ -337,8 +404,14 @@ class hipGraphKernelNode : public hipGraphNode {
void GetParams(hipKernelNodeParams* params) {
std::memcpy(params, pKernelParams_, sizeof(hipKernelNodeParams));
}
void SetParams(const hipKernelNodeParams* params) {
hipError_t SetParams(const hipKernelNodeParams* params) {
// updates kernel params
hipError_t status = ihipValidateKernelParams(params);
if (hipSuccess != status) {
return status;
}
std::memcpy(pKernelParams_, params, sizeof(hipKernelNodeParams));
return hipSuccess;
}
hipError_t SetCommandParams(const hipKernelNodeParams* params) {
if (params->func != pKernelParams_->func) {
@@ -378,9 +451,13 @@ class hipGraphMemcpyNode : public hipGraphNode {
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
commands_.reserve(1);
amd::Command* command;
hipError_t status = ihipMemcpy3DCommand(command, pCopyParams_, queue);
status = ihipMemcpy3DCommand(command, pCopyParams_, queue);
commands_.emplace_back(command);
return status;
}
@@ -388,10 +465,16 @@ class hipGraphMemcpyNode : public hipGraphNode {
void GetParams(hipMemcpy3DParms* params) {
std::memcpy(params, pCopyParams_, sizeof(hipMemcpy3DParms));
}
void SetParams(const hipMemcpy3DParms* params) {
hipError_t SetParams(const hipMemcpy3DParms* params) {
hipError_t status = ValidateParams(params);
if (status != hipSuccess) {
return status;
}
std::memcpy(pCopyParams_, params, sizeof(hipMemcpy3DParms));
return hipSuccess;
}
hipError_t SetCommandParams(const hipMemcpy3DParms* pNodeParams);
hipError_t ValidateParams(const hipMemcpy3DParms* pNodeParams);
};
class hipGraphMemcpyNode1D : public hipGraphNode {
@@ -411,23 +494,41 @@ class hipGraphMemcpyNode1D : public hipGraphNode {
hipGraphNode* clone() const {
return new hipGraphMemcpyNode1D(static_cast<hipGraphMemcpyNode1D const&>(*this));
}
virtual hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
commands_.reserve(1);
amd::Command* command = nullptr;
hipError_t status = ihipMemcpyCommand(command, dst_, src_, count_, kind_, *queue);
status = ihipMemcpyCommand(command, dst_, src_, count_, kind_, *queue);
commands_.emplace_back(command);
return status;
}
void SetParams(void* dst, const void* src, size_t count, hipMemcpyKind kind) {
void EnqueueCommands(hipStream_t stream) {
if (!commands_.empty()) {
for (auto& command : commands_) {
command->enqueue();
}
}
}
hipError_t SetParams(void* dst, const void* src, size_t count, hipMemcpyKind kind) {
hipError_t status = ValidateParams(dst, src, count, kind);
if (status != hipSuccess) {
return status;
}
dst_ = dst;
src_ = src;
count_ = count;
kind_ = kind;
return hipSuccess;
}
hipError_t SetCommandParams(void* dst, const void* src, size_t count, hipMemcpyKind kind);
hipError_t ValidateParams(void* dst, const void* src, size_t count, hipMemcpyKind kind);
};
class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D {
@@ -449,12 +550,16 @@ class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D {
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
commands_.reserve(1);
amd::Command* command = nullptr;
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipError_t status = ihipMemcpySymbol_validate(symbol_, count_, offset_, sym_size, device_ptr);
status = ihipMemcpySymbol_validate(symbol_, count_, offset_, sym_size, device_ptr);
if (status != hipSuccess) {
return status;
}
@@ -466,12 +571,21 @@ class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D {
return status;
}
void SetParams(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind) {
hipError_t SetParams(void* dst, const void* symbol, size_t count, size_t offset,
hipMemcpyKind kind) {
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipError_t status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
if (status != hipSuccess) {
return status;
}
dst_ = dst;
symbol_ = symbol;
count_ = count;
offset_ = offset;
kind_ = kind;
return hipSuccess;
}
hipError_t SetCommandParams(void* dst, const void* symbol, size_t count, size_t offset,
@@ -504,12 +618,16 @@ class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D {
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
commands_.reserve(1);
amd::Command* command = nullptr;
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipError_t status = ihipMemcpySymbol_validate(symbol_, count_, offset_, sym_size, device_ptr);
status = ihipMemcpySymbol_validate(symbol_, count_, offset_, sym_size, device_ptr);
if (status != hipSuccess) {
return status;
}
@@ -521,13 +639,21 @@ class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D {
return status;
}
void SetParams(const void* symbol, const void* src, size_t count, size_t offset,
hipMemcpyKind kind) {
hipError_t SetParams(const void* symbol, const void* src, size_t count, size_t offset,
hipMemcpyKind kind) {
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipError_t status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
if (status != hipSuccess) {
return status;
}
symbol_ = symbol;
src_ = src;
count_ = count;
offset_ = offset;
kind_ = kind;
return hipSuccess;
}
hipError_t SetCommandParams(const void* symbol, const void* src, size_t count, size_t offset,
@@ -561,6 +687,10 @@ class hipGraphMemsetNode : public hipGraphNode {
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
if (pMemsetParams_->height == 1) {
return ihipMemsetCommand(commands_, pMemsetParams_->dst, pMemsetParams_->value,
pMemsetParams_->elementSize,
@@ -578,8 +708,15 @@ class hipGraphMemsetNode : public hipGraphNode {
void GetParams(hipMemsetParams* params) {
std::memcpy(params, pMemsetParams_, sizeof(hipMemsetParams));
}
void SetParams(const hipMemsetParams* params) {
hipError_t SetParams(const hipMemsetParams* params) {
hipError_t hip_error = hipSuccess;
hip_error = ihipMemset_validate(params->dst, params->value, params->elementSize,
params->width * params->elementSize);
if (hip_error != hipSuccess) {
return hip_error;
}
std::memcpy(pMemsetParams_, params, sizeof(hipMemsetParams));
return hipSuccess;
}
};
@@ -596,10 +733,14 @@ class hipGraphEventRecordNode : public hipGraphNode {
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
hip::Event* e = reinterpret_cast<hip::Event*>(event_);
commands_.reserve(1);
amd::Command* command;
hipError_t status = e->recordCommand(command, queue);
status = e->recordCommand(command, queue);
commands_.emplace_back(command);
return status;
}
@@ -618,16 +759,9 @@ class hipGraphEventRecordNode : public hipGraphNode {
void GetParams(hipEvent_t* event) { *event = event_; }
void SetParams(hipEvent_t event) { event_ = event; }
hipError_t SetExecParams(hipEvent_t event) {
amd::HostQueue* queue;
if (!commands_.empty()) {
queue = commands_[0]->queue();
commands_[0]->release();
}
commands_.clear();
return CreateCommand(queue);
hipError_t SetParams(hipEvent_t event) {
event_ = event;
return hipSuccess;
}
};
@@ -644,10 +778,14 @@ class hipGraphEventWaitNode : public hipGraphNode {
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
hip::Event* e = reinterpret_cast<hip::Event*>(event_);
commands_.reserve(1);
amd::Command* command;
hipError_t status = e->streamWaitCommand(command, queue);
status = e->streamWaitCommand(command, queue);
commands_.emplace_back(command);
return status;
}
@@ -666,16 +804,9 @@ class hipGraphEventWaitNode : public hipGraphNode {
void GetParams(hipEvent_t* event) { *event = event_; }
void SetParams(hipEvent_t event) { event_ = event; }
hipError_t SetExecParams(hipEvent_t event) {
amd::HostQueue* queue;
if (!commands_.empty()) {
queue = commands_[0]->queue();
commands_[0]->release();
}
commands_.clear();
return CreateCommand(queue);
hipError_t SetParams(hipEvent_t event) {
event_ = event;
return hipSuccess;
}
};
@@ -701,8 +832,9 @@ class hipGraphHostNode : public hipGraphNode {
void GetParams(hipHostNodeParams* params) {
std::memcpy(params, pNodeParams_, sizeof(hipHostNodeParams));
}
void SetParams(hipHostNodeParams* params) {
hipError_t SetParams(hipHostNodeParams* params) {
std::memcpy(pNodeParams_, params, sizeof(hipHostNodeParams));
return hipSuccess;
}
};
@@ -716,6 +848,10 @@ class hipGraphEmptyNode : public hipGraphNode {
}
hipError_t CreateCommand(amd::HostQueue* queue) {
hipError_t status = hipGraphNode::CreateCommand(queue);
if (status != hipSuccess) {
return status;
}
amd::Command::EventWaitList waitList;
commands_.reserve(1);
amd::Command* command = new amd::Marker(*queue, !kMarkerDisableFlush, waitList);
@@ -723,43 +859,3 @@ class hipGraphEmptyNode : public hipGraphNode {
return hipSuccess;
}
};
struct hipGraphExec {
std::vector<std::vector<Node>> parallelLists_;
// level order of the graph doesn't include nodes embedded as part of the child graph
std::vector<Node> levelOrder_;
std::unordered_map<Node, std::vector<Node>> nodeWaitLists_;
std::vector<amd::HostQueue*> parallelQueues_;
static std::unordered_map<amd::Command*, hipGraphExec_t> activeGraphExec_;
amd::Command::EventWaitList graphLastCmdWaitList_;
amd::Command* lastEnqueuedGraphCmd_;
std::atomic<bool> bExecPending_;
amd::Command* rootCommand_;
public:
hipGraphExec(std::vector<Node>& levelOrder, std::vector<std::vector<Node>>& lists,
std::unordered_map<Node, std::vector<Node>>& nodeWaitLists)
: parallelLists_(lists),
levelOrder_(levelOrder),
nodeWaitLists_(nodeWaitLists),
lastEnqueuedGraphCmd_(nullptr),
rootCommand_(nullptr) {
bExecPending_.store(false);
}
~hipGraphExec() {
for (auto queue : parallelQueues_) {
queue->release();
}
for (auto node : levelOrder_) {
node->ReleaseCommand();
}
}
hipError_t CreateQueues();
hipError_t FillCommands();
hipError_t Init();
void UpdateGraphToWaitOnRoot();
hipError_t Run(hipStream_t stream);
static void ResetGraph(cl_event event, cl_int command_exec_status, void* user_data);
};