From abb6f921380d48953004d8bbf9b69920edadd788 Mon Sep 17 00:00:00 2001 From: anusha GodavarthySurya Date: Wed, 17 Nov 2021 00:13:52 -0800 Subject: [PATCH] SWDEV-240806 - hipGraph performance create new graph commands for every launch Change-Id: I7360b48055f67a91c23169ca532a19124d40e98b [ROCm/clr commit: ec6b8d82b52cc4a4c1e951a66f5cacc330914721] --- projects/clr/hipamd/src/hip_graph.cpp | 85 ++-- .../clr/hipamd/src/hip_graph_internal.cpp | 385 +++++++++++++----- .../clr/hipamd/src/hip_graph_internal.hpp | 376 ++++++++++------- 3 files changed, 564 insertions(+), 282 deletions(-) diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 3c3257536c..30c87e3d2e 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -25,7 +25,6 @@ #include "hip_event.hpp" thread_local std::vector g_captureStreams; -std::unordered_map 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(node)->SetParams(dst, src, count, kind); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(node)->SetCommandParams(dst, src, count, kind)); + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(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 clonedNodes; + hipGraph_t clonedGraph = graph->clone(clonedNodes); std::vector> parallelLists; std::unordered_map> nodeWaitLists; - graph->GetRunList(parallelLists, nodeWaitLists); + clonedGraph->GetRunList(parallelLists, nodeWaitLists); std::vector 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(node)->SetParams(pNodeParams); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(node)->SetParams(pNodeParams); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(node)->SetCommandParams(pNodeParams)); + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(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(node)->SetParams(pNodeParams); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(node)->SetCommandParams(pNodeParams)); + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(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(node)->SetParams(dst, symbol, count, offset, - kind); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(node)->SetCommandParams( - dst, symbol, count, offset, kind)); + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(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(node)->SetParams(symbol, src, count, offset, kind); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(node)->SetCommandParams( - symbol, src, count, offset, kind)); + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(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(node)->SetParams(event); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(hNode)->SetExecParams(event)); + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(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(node)->SetParams(event); - HIP_RETURN(hipSuccess); + HIP_RETURN(reinterpret_cast(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(hNode)->SetExecParams(event)); + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(event)); } diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index 89316b4c7d..8d465145ab 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -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(&pCopy)->srcHost = pCopy.srcDevice; + } + } + hipMemoryType dstMemoryType = pCopy.dstMemoryType; + if (dstMemoryType == hipMemoryTypeUnified) { + dstMemoryType = + amd::MemObjMap::FindMemObj(pCopy.dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + if (srcMemoryType == hipMemoryTypeHost) { + const_cast(&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(&pCopy)->srcDevice = const_cast(pCopy.srcHost); + } + } + if (dstMemoryType == hipMemoryTypeHost) { + amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.dstHost); + dstMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; + if (dstMemoryType == hipMemoryTypeDevice) { + const_cast(&pCopy)->dstDevice = const_cast(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& levelOrder) { } } -ihipGraph* ihipGraph::clone() const { +ihipGraph* ihipGraph::clone(std::unordered_map& clonedNodes) const { ihipGraph* newGraph = new ihipGraph(); - std::unordered_map 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 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>& parallelLists, + std::unordered_map>& nodeWaitLists, + std::vector& 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(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>& 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; } diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 002c13f1b0..360356c25a 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -34,8 +34,15 @@ typedef hipGraphNode* Node; hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams); +hipError_t FillCommands(std::vector>& parallelLists, + std::unordered_map>& nodeWaitLists, + std::vector& levelOrder, amd::Command*& rootCommand, + amd::Command*& endCommand, amd::HostQueue* queue); +void UpdateQueue(std::vector>& 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& GetCommands() { return commands_; } + virtual std::vector& 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& 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>& parallelList, std::unordered_map>& dependencies); void LevelOrder(std::vector& levelOrder); + ihipGraph* clone(std::unordered_map& clonedNodes) const; ihipGraph* clone() const; }; +struct hipGraphExec { + std::vector> parallelLists_; + // level order of the graph doesn't include nodes embedded as part of the child graph + std::vector levelOrder_; + std::unordered_map> nodeWaitLists_; + std::vector parallelQueues_; + uint currentQueueIndex_; + std::unordered_map clonedNodes_; + amd::Command* lastEnqueuedCommand_; + + public: + hipGraphExec(std::vector& levelOrder, std::vector>& lists, + std::unordered_map>& nodeWaitLists, + std::unordered_map& 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 childGraphlevelOrder_; + std::vector> parallelLists_; + std::unordered_map> 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 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 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& 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>& parallelList, std::unordered_map>& dependencies) { - childGraph_->GetRunList(parallelList, dependencies); + childGraph_->GetRunList(parallelLists_, nodeWaitLists_); } void LevelOrder(std::vector& 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(*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(*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(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(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> parallelLists_; - // level order of the graph doesn't include nodes embedded as part of the child graph - std::vector levelOrder_; - std::unordered_map> nodeWaitLists_; - std::vector parallelQueues_; - static std::unordered_map activeGraphExec_; - amd::Command::EventWaitList graphLastCmdWaitList_; - amd::Command* lastEnqueuedGraphCmd_; - std::atomic bExecPending_; - amd::Command* rootCommand_; - - public: - hipGraphExec(std::vector& levelOrder, std::vector>& lists, - std::unordered_map>& 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); -};