diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index f5efa51032..a5b3c556ae 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -974,6 +974,10 @@ typedef enum cudaStreamCaptureStatus hipStreamCaptureStatus; #define hipStreamCaptureStatusActive cudaStreamCaptureStatusActive #define hipStreamCaptureStatusInvalidated cudaStreamCaptureStatusInvalidated +typedef enum cudaStreamUpdateCaptureDependenciesFlags hipStreamUpdateCaptureDependenciesFlags; +#define hipStreamAddCaptureDependencies cudaStreamAddCaptureDependencies +#define hipStreamSetCaptureDependencies cudaStreamSetCaptureDependencies + /** * Stream CallBack struct */ @@ -2373,48 +2377,137 @@ inline static hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGra cudaGraphAddEmptyNode(pGraphNode, graph, pDependencies, numDependencies)); } -inline static hipError_t hipStreamWriteValue32(hipStream_t stream, - void* ptr, int32_t value, unsigned int flags) { - if (value < 0) { - printf("Warning! value is negative, CUDA accept positive values\n"); - } - return hipCUResultTohipError(cuStreamWriteValue32(stream, reinterpret_cast(ptr), - static_cast(value), flags)); +inline static hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, int32_t value, + unsigned int flags) { + if (value < 0) { + printf("Warning! value is negative, CUDA accept positive values\n"); + } + return hipCUResultTohipError(cuStreamWriteValue32(stream, reinterpret_cast(ptr), + static_cast(value), flags)); } -inline static hipError_t hipStreamWriteValue64(hipStream_t stream, - void* ptr, int64_t value, unsigned int flags) { - if (value < 0) { - printf("Warning! value is negative, CUDA accept positive values\n"); - } - return hipCUResultTohipError(cuStreamWriteValue64(stream, reinterpret_cast(ptr), - static_cast(value), flags)); +inline static hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, int64_t value, + unsigned int flags) { + if (value < 0) { + printf("Warning! value is negative, CUDA accept positive values\n"); + } + return hipCUResultTohipError(cuStreamWriteValue64(stream, reinterpret_cast(ptr), + static_cast(value), flags)); } -inline static hipError_t hipStreamWaitValue32(hipStream_t stream, - void* ptr, int32_t value, unsigned int flags, +inline static hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, int32_t value, + unsigned int flags, uint32_t mask __dparm(0xFFFFFFFF)) { - if (value < 0) { - printf("Warning! value is negative, CUDA accept positive values\n"); - } - if (mask != STREAM_OPS_WAIT_MASK_32) { - printf("Warning! mask will not have impact as CUDA ignores it.\n"); - } - return hipCUResultTohipError(cuStreamWaitValue32(stream, reinterpret_cast(ptr), - static_cast(value), flags)); + if (value < 0) { + printf("Warning! value is negative, CUDA accept positive values\n"); + } + if (mask != STREAM_OPS_WAIT_MASK_32) { + printf("Warning! mask will not have impact as CUDA ignores it.\n"); + } + return hipCUResultTohipError(cuStreamWaitValue32(stream, reinterpret_cast(ptr), + static_cast(value), flags)); } -inline static hipError_t hipStreamWaitValue64(hipStream_t stream, - void* ptr, int64_t value, unsigned int flags, +inline static hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, int64_t value, + unsigned int flags, uint64_t mask __dparm(0xFFFFFFFFFFFFFFFF)) { - if (value < 0) { - printf("Warning! value is negative, CUDA accept positive values\n"); - } - if (mask != STREAM_OPS_WAIT_MASK_64) { - printf("Warning! mask will not have impact as CUDA ignores it.\n"); - } - return hipCUResultTohipError(cuStreamWaitValue64(stream, reinterpret_cast(ptr), - static_cast(value), flags)); + if (value < 0) { + printf("Warning! value is negative, CUDA accept positive values\n"); + } + if (mask != STREAM_OPS_WAIT_MASK_64) { + printf("Warning! mask will not have impact as CUDA ignores it.\n"); + } + return hipCUResultTohipError(cuStreamWaitValue64(stream, reinterpret_cast(ptr), + static_cast(value), flags)); +} + +inline static hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from, + const hipGraphNode_t* to, + size_t numDependencies) { + return hipCUDAErrorTohipError(cudaGraphRemoveDependencies(graph, from, to, numDependencies)); +} + + +inline static hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, + hipGraphNode_t* to, size_t* numEdges) { + return hipCUDAErrorTohipError(cudaGraphGetEdges(graph, from, to, numEdges)); +} + + +inline static hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, + hipGraphNode_t* pDependencies, + size_t* pNumDependencies) { + return hipCUDAErrorTohipError( + cudaGraphNodeGetDependencies(node, pDependencies, pNumDependencies)); +} + +inline static hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, + hipGraphNode_t* pDependentNodes, + size_t* pNumDependentNodes) { + return hipCUDAErrorTohipError( + cudaGraphNodeGetDependentNodes(node, pDependentNodes, pNumDependentNodes)); +} + +inline static hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType) { + return hipCUDAErrorTohipError(cudaGraphNodeGetType(node, pType)); +} + +inline static hipError_t hipGraphDestroyNode(hipGraphNode_t node) { + return hipCUDAErrorTohipError(cudaGraphDestroyNode(node)); +} + +inline static hipError_t hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph) { + return hipCUDAErrorTohipError(cudaGraphClone(pGraphClone, originalGraph)); +} + +inline static hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t originalNode, + hipGraph_t clonedGraph) { + return hipCUDAErrorTohipError(cudaGraphNodeFindInClone(pNode, originalNode, clonedGraph)); +} + +inline static hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, + size_t numDependencies, hipGraph_t childGraph) { + return hipCUDAErrorTohipError( + cudaGraphAddChildGraphNode(pGraphNode, graph, pDependencies, numDependencies, childGraph)); +} + +inline static hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph) { + return hipCUDAErrorTohipError(cudaGraphChildGraphNodeGetGraph(node, pGraph)); +} + +inline static hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, + hipGraphNode_t node, + hipGraph_t childGraph) { + return hipCUDAErrorTohipError( + cudaGraphExecChildGraphNodeSetParams(hGraphExec, node, childGraph)); +} + + +inline static hipError_t hipStreamGetCaptureInfo(hipStream_t stream, + hipStreamCaptureStatus* pCaptureStatus, + unsigned long long* pId) { + return hipCUDAErrorTohipError(cudaStreamGetCaptureInfo(stream, pCaptureStatus, pId)); +} + +inline static hipError_t hipStreamGetCaptureInfo_v2( + hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, + unsigned long long* id_out __dparm(0), hipGraph_t* graph_out __dparm(0), + const hipGraphNode_t** dependencies_out __dparm(0), size_t* numDependencies_out __dparm(0)) { + return hipCUDAErrorTohipError(cudaStreamGetCaptureInfo_v2( + stream, captureStatus_out, id_out, graph_out, dependencies_out, numDependencies_out)); +} + +inline static hipError_t hipStreamIsCapturing(hipStream_t stream, + hipStreamCaptureStatus* pCaptureStatus) { + return hipCUDAErrorTohipError(cudaStreamIsCapturing(stream, pCaptureStatus)); +} + +inline static hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, + hipGraphNode_t* dependencies, + size_t numDependencies, + unsigned int flags __dparm(0)) { + return hipCUDAErrorTohipError(cudaStreamUpdateCaptureDependencies(stream, dependencies, flags)); } #endif //__CUDACC__ diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index f6de32b484..5864df5c27 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -27,6 +27,14 @@ 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) { + graph->AddNode(graphNode); + for (size_t i = 0; i < numDependencies; i++) { + pDependencies[i]->AddEdge(graphNode); + } +} + hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams) { hipFunction_t func = nullptr; hipError_t status = @@ -70,14 +78,7 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, return hipErrorInvalidDeviceFunction; } *pGraphNode = new hipGraphKernelNode(pNodeParams, func); - if (numDependencies == 0) { - graph->AddNode(*pGraphNode); - } - for (size_t i = 0; i < numDependencies; i++) { - if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { - return hipErrorInvalidValue; - } - } + ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); return hipSuccess; } @@ -88,16 +89,12 @@ hipError_t ihipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, (numDependencies > 0 && pDependencies == nullptr) || pCopyParams == nullptr) { return hipErrorInvalidValue; } - ihipMemcpy3D_validate(pCopyParams); + hipError_t status = ihipMemcpy3D_validate(pCopyParams); + if (status != hipSuccess) { + return status; + } *pGraphNode = new hipGraphMemcpyNode(pCopyParams); - if (numDependencies == 0) { - graph->AddNode(*pGraphNode); - } - for (size_t i = 0; i < numDependencies; i++) { - if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { - return hipErrorInvalidValue; - } - } + ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); return hipSuccess; } @@ -108,16 +105,12 @@ hipError_t ihipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph (numDependencies > 0 && pDependencies == nullptr)) { return hipErrorInvalidValue; } - ihipMemcpy_validate(dst, src, count, kind); + hipError_t status = ihipMemcpy_validate(dst, src, count, kind); + if (status != hipSuccess) { + return status; + } *pGraphNode = new hipGraphMemcpyNode1D(dst, src, count, kind); - if (numDependencies == 0) { - graph->AddNode(*pGraphNode); - } - for (size_t i = 0; i < numDependencies; i++) { - if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { - return hipErrorInvalidValue; - } - } + ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); return hipSuccess; } @@ -128,25 +121,22 @@ hipError_t ihipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, (numDependencies > 0 && pDependencies == nullptr) || pMemsetParams == nullptr) { return hipErrorInvalidValue; } + hipError_t status; if (pMemsetParams->height == 1) { - ihipMemset_validate(pMemsetParams->dst, pMemsetParams->value, pMemsetParams->elementSize, - pMemsetParams->width * pMemsetParams->elementSize); + status = + ihipMemset_validate(pMemsetParams->dst, pMemsetParams->value, pMemsetParams->elementSize, + pMemsetParams->width * pMemsetParams->elementSize); } else { auto sizeBytes = pMemsetParams->width * pMemsetParams->height * 1; - ihipMemset3D_validate( + status = ihipMemset3D_validate( {pMemsetParams->dst, pMemsetParams->pitch, pMemsetParams->width, pMemsetParams->height}, pMemsetParams->value, {pMemsetParams->width, pMemsetParams->height, 1}, sizeBytes); } - + if (status != hipSuccess) { + return status; + } *pGraphNode = new hipGraphMemsetNode(pMemsetParams); - if (numDependencies == 0) { - graph->AddNode(*pGraphNode); - } - for (size_t i = 0; i < numDependencies; i++) { - if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { - return hipErrorInvalidValue; - } - } + ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); return hipSuccess; } @@ -194,22 +184,17 @@ hipError_t capturehipMemcpyAsync(hipStream_t& stream, void*& dst, const void*& s ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy1D on stream : %p", stream); hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; hipGraph_t graph = nullptr; std::vector pDependencies = s->GetLastCapturedNodes(); size_t numDependencies = s->GetLastCapturedNodes().size(); graph = s->GetCaptureGraph(); - ihipMemcpy_validate(dst, src, sizeBytes, kind); - pGraphNode = new hipGraphMemcpyNode1D(dst, src, sizeBytes, kind); - if (numDependencies == 0) { - graph->AddNode(pGraphNode); + hipError_t status = ihipMemcpy_validate(dst, src, sizeBytes, kind); + if (status != hipSuccess) { + return status; } - for (size_t i = 0; i < numDependencies; i++) { - if (graph->AddEdge(pDependencies[i], pGraphNode) != hipSuccess) { - return hipErrorInvalidValue; - } - } - s->SetLastCapturedNode(pGraphNode); + hipGraphNode_t node = new hipGraphMemcpyNode1D(dst, src, sizeBytes, kind); + ihipGraphAddNode(node, graph, pDependencies.data(), numDependencies); + s->SetLastCapturedNode(node); return hipSuccess; } @@ -318,13 +303,12 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns HIP_RETURN(hipSuccess); } -hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus** pCaptureStatus) { +hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) { HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus); if (stream == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - hipStreamCaptureStatus captureStatus = reinterpret_cast(stream)->GetCaptureStatus(); - *pCaptureStatus = &captureStatus; + *pCaptureStatus = reinterpret_cast(stream)->GetCaptureStatus(); HIP_RETURN(hipSuccess); } @@ -363,8 +347,22 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { HIP_RETURN(hipErrorStreamCaptureInvalidated); } // check if all parallel streams have joined + // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do + // not result in hipErrorStreamCaptureUnjoined if (s->GetCaptureGraph()->GetLeafNodeCount() != 1) { - return hipErrorStreamCaptureUnjoined; + std::vector leafNodes = s->GetCaptureGraph()->GetLeafNodes(); + const std::vector& removedDepNodes = s->GetRemovedDependencies(); + bool foundInRemovedDep = false; + for (auto leafNode : leafNodes) { + for (auto node : removedDepNodes) { + if (node == leafNode) { + foundInRemovedDep = true; + } + } + } + if (foundInRemovedDep == false) { + return hipErrorStreamCaptureUnjoined; + } } *pGraph = s->GetCaptureGraph(); // end capture on all streams/events part of graph capture @@ -379,6 +377,9 @@ hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags) { hipError_t hipGraphDestroy(hipGraph_t graph) { HIP_INIT_API(hipGraphDestroy, graph); + if (graph == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } delete graph; HIP_RETURN(hipSuccess); } @@ -430,15 +431,21 @@ hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, HIP_RETURN(hipErrorInvalidValue); } *pGraphNode = new hipGraphEmptyNode(); - if (numDependencies == 0) { - graph->AddNode(*pGraphNode); + ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + hipGraph_t childGraph) { + HIP_INIT_API(hipGraphAddChildGraphNode, pGraphNode, pDependencies, numDependencies, childGraph); + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr) || childGraph == nullptr) { + HIP_RETURN(hipErrorInvalidValue); } - for (size_t i = 0; i < numDependencies; i++) { - if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { - return hipErrorInvalidValue; - } - } - return hipSuccess; + *pGraphNode = new hipChildGraphNode(childGraph); + ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); + HIP_RETURN(hipSuccess); } hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, @@ -464,6 +471,9 @@ hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) { HIP_INIT_API(hipGraphExecDestroy, pGraphExec); + if (pGraphExec == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } delete pGraphExec; HIP_RETURN(hipSuccess); } @@ -474,6 +484,9 @@ 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) { + HIP_RETURN(hipErrorInvalidValue); + } HIP_RETURN_DURATION(ihipGraphlaunch(graphExec, stream)); } @@ -482,22 +495,43 @@ hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* num if (graph == nullptr || numNodes == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *numNodes = graph->GetNodeCount(); - if (*numNodes > 0) { - nodes = graph->GetNodes().data(); + const std::vector& graphNodes = graph->GetNodes(); + if (nodes == nullptr) { + *numNodes = graphNodes.size(); + HIP_RETURN(hipSuccess); + } else if (*numNodes < graphNodes.size()) { + HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(hipSuccess); + for (int i = 0; i < graphNodes.size(); i++) { + nodes[i] = graphNodes[i]; + } + for (int i = graphNodes.size(); i < *numNodes; i++) { + nodes[i] = nullptr; + } + *numNodes = graphNodes.size(); } hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, size_t* pNumRootNodes) { HIP_INIT_API(hipGraphGetRootNodes, graph, pRootNodes, pNumRootNodes); + if (graph == nullptr || pNumRootNodes == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - std::vector rootNodes = graph->GetRootNodes(); - pRootNodes = rootNodes.data(); - *pNumRootNodes = rootNodes.size(); + const std::vector nodes = graph->GetRootNodes(); + if (pRootNodes == nullptr) { + *pNumRootNodes = nodes.size(); + HIP_RETURN(hipSuccess); + } else if (*pNumRootNodes < nodes.size()) { + HIP_RETURN(hipErrorInvalidValue); + } + for (int i = 0; i < nodes.size(); i++) { + pRootNodes[i] = nodes[i]; + } + for (int i = nodes.size(); i < *pNumRootNodes; i++) { + pRootNodes[i] = nullptr; + } + *pNumRootNodes = nodes.size(); HIP_RETURN(hipSuccess); } @@ -568,9 +602,7 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, HIP_RETURN(hipErrorInvalidValue); } for (size_t i = 0; i < numDependencies; i++) { - if (graph->AddEdge(from[i], to[i]) != hipSuccess) { - HIP_RETURN(hipErrorInvalidValue); - } + from[i]->AddEdge(to[i]); } HIP_RETURN(hipSuccess); } @@ -583,3 +615,220 @@ hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo } return reinterpret_cast(node)->SetCommandParams(pNodeParams); } + +hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph) { + HIP_INIT_API(hipGraphChildGraphNodeGetGraph, node, pGraph); + if (node == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + *pGraph = reinterpret_cast(node)->GetChildGraph(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + hipGraph_t childGraph) { + HIP_INIT_API(hipGraphExecChildGraphNodeSetParams, hGraphExec, node, childGraph); + if (hGraphExec == nullptr || node == nullptr || childGraph == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(hipSuccess); +} + +hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus, + unsigned long long* pId) { + HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId); + if (pCaptureStatus == nullptr || pId == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + if (stream == nullptr) { + HIP_RETURN(hipErrorStreamCaptureImplicit); + } + hip::Stream* s = reinterpret_cast(stream); + *pCaptureStatus = s->GetCaptureStatus(); + *pId = s->GetCaptureID(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, + unsigned long long* id_out, hipGraph_t* graph_out, + const hipGraphNode_t** dependencies_out, + size_t* numDependencies_out) { + HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out, + dependencies_out, numDependencies_out); + if (stream == nullptr) { + HIP_RETURN(hipErrorStreamCaptureImplicit); + } + hip::Stream* s = reinterpret_cast(stream); + *captureStatus_out = s->GetCaptureStatus(); + if (*captureStatus_out == hipStreamCaptureStatusActive) { + if (id_out != nullptr) { + *id_out = s->GetCaptureID(); + } + if (graph_out != nullptr) { + *graph_out = s->GetCaptureGraph(); + } + if (dependencies_out != nullptr && numDependencies_out != nullptr) { + *dependencies_out = s->GetLastCapturedNodes().data(); + *numDependencies_out = s->GetLastCapturedNodes().size(); + } + } + HIP_RETURN(hipSuccess); +} + +hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, + size_t numDependencies, unsigned int flags) { + HIP_INIT_API(hipStreamUpdateCaptureDependencies, stream, dependencies, numDependencies, flags); + hip::Stream* s = reinterpret_cast(stream); + if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) { + return hipErrorIllegalState; + } + if (numDependencies > 0 && dependencies == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + std::vector depNodes; + for (int i = 0; i < numDependencies; i++) { + depNodes.push_back(dependencies[i]); + } + if (flags == hipStreamAddCaptureDependencies) { + s->AddCrossCapturedNode(depNodes); + } else if (flags == hipStreamSetCaptureDependencies) { + bool replace = true; + s->AddCrossCapturedNode(depNodes, replace); + } + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from, + const hipGraphNode_t* to, size_t numDependencies) { + HIP_INIT_API(hipGraphRemoveDependencies, graph, from, to, numDependencies); + if (graph == nullptr || from == nullptr || to == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + for (size_t i = 0; i < numDependencies; i++) { + from[i]->RemoveEdge(to[i]); + } + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode_t* to, + size_t* numEdges) { + HIP_INIT_API(hipGraphGetEdges, graph, from, to, numEdges); + if (graph == nullptr || numEdges == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + const std::vector> edges = graph->GetEdges(); + // returns only the number of edges in numEdges when from and to are null + if (from == nullptr && to == nullptr) { + *numEdges = edges.size(); + return hipSuccess; + } else if (*numEdges < edges.size()) { + HIP_RETURN(hipErrorInvalidValue); + } + for (int i = 0; i < edges.size(); i++) { + from[i] = edges[i].first; + to[i] = edges[i].second; + } + // If numEdges > actual number of edges, the remaining entries in from and to will be set to NULL + for (int i = edges.size(); i < *numEdges; i++) { + from[i] = nullptr; + to[i] = nullptr; + } + *numEdges = edges.size(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDependencies, + size_t* pNumDependencies) { + HIP_INIT_API(hipGraphNodeGetDependencies, node, pDependencies, pNumDependencies); + if (node == nullptr || pNumDependencies == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + const std::vector& dependencies = node->GetDependencies(); + if (pDependencies == NULL) { + *pNumDependencies = dependencies.size(); + HIP_RETURN(hipSuccess); + } else if (*pNumDependencies < dependencies.size()) { + HIP_RETURN(hipErrorInvalidValue); + } + + for (int i = 0; i < dependencies.size(); i++) { + pDependencies[i] = dependencies[i]; + } + // pNumDependencies > actual number of dependencies, the remaining entries in pDependencies will + // be set to NULL + for (int i = dependencies.size(); i < *pNumDependencies; i++) { + pDependencies[i] = nullptr; + } + *pNumDependencies = dependencies.size(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pDependentNodes, + size_t* pNumDependentNodes) { + HIP_INIT_API(hipGraphNodeGetDependentNodes, node, pDependentNodes, pNumDependentNodes); + if (node == nullptr || pNumDependentNodes == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + const std::vector& dependents = node->GetEdges(); + if (pDependentNodes == NULL) { + *pNumDependentNodes = dependents.size(); + HIP_RETURN(hipSuccess); + } else if (*pNumDependentNodes < dependents.size()) { + HIP_RETURN(hipErrorInvalidValue); + } + + for (int i = 0; i < dependents.size(); i++) { + pDependentNodes[i] = dependents[i]; + } + // pNumDependentNodes > actual number of dependents, the remaining entries in pDependentNodes will + // be set to NULL + for (int i = dependents.size(); i < *pNumDependentNodes; i++) { + pDependentNodes[i] = nullptr; + } + *pNumDependentNodes = dependents.size(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType) { + HIP_INIT_API(hipGraphNodeGetType, node, pType); + if (node == nullptr || pType == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + *pType = node->GetType(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphDestroyNode(hipGraphNode_t node) { + HIP_INIT_API(hipGraphDestroyNode, node); + if (node == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + node->GetParentGraph()->RemoveNode(node); + // Takescare of removing its dependencies and dependent nodes + delete node; + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph) { + HIP_INIT_API(hipGraphClone, pGraphClone, originalGraph); + if (originalGraph == nullptr || pGraphClone == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + *pGraphClone = originalGraph->clone(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t originalNode, + hipGraph_t clonedGraph) { + HIP_INIT_API(hipGraphNodeFindInClone, pNode, originalNode, clonedGraph); + if (pNode == nullptr || originalNode == nullptr || clonedGraph == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + for (auto node : clonedGraph->GetNodes()) { + if (node->GetID() == originalNode->GetID()) { + *pNode = node; + HIP_RETURN(hipSuccess); + } + } + HIP_RETURN(hipErrorInvalidValue); +} diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index fe6618d8db..b414478e85 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -45,49 +45,24 @@ const char* GetGraphNodeTypeString(uint32_t op) { return case_string; }; -hipError_t ihipGraph::AddNode(const Node& node) { +int hipGraphNode::nextID = 0; + +void ihipGraph::AddNode(const Node& node) { vertices_.emplace_back(node); - nodeOutDegree_[node] = 0; - nodeInDegree_[node] = 0; - node->SetLevel(0); ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add %s(%p)\n", GetGraphNodeTypeString(node->GetType()), node); - return hipSuccess; + node->SetParentGraph(this); } -hipError_t ihipGraph::AddEdge(const Node& parentNode, const Node& childNode) { - // if vertice doesn't exist, add it to the graph - if (std::find(vertices_.begin(), vertices_.end(), parentNode) == vertices_.end()) { - AddNode(parentNode); - } - if (std::find(vertices_.begin(), vertices_.end(), childNode) == vertices_.end()) { - AddNode(childNode); - } - // Check if edge already exists - auto connectedEdges = edges_.find(parentNode); - if (connectedEdges != edges_.end()) { - if (std::find(connectedEdges->second.begin(), connectedEdges->second.end(), childNode) != - connectedEdges->second.end()) { - return hipSuccess; - } - connectedEdges->second.emplace_back(childNode); - } else { - edges_[parentNode] = {childNode}; - } - nodeOutDegree_[parentNode]++; - nodeInDegree_[childNode]++; - childNode->SetLevel(std::max(childNode->GetLevel(), parentNode->GetLevel() + 1)); - ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add edge btwn %s(%p) - %s(%p)\n", - GetGraphNodeTypeString(parentNode->GetType()), parentNode, - GetGraphNodeTypeString(childNode->GetType()), childNode); - return hipSuccess; +void ihipGraph::RemoveNode(const Node& node) { + vertices_.erase(std::remove(vertices_.begin(), vertices_.end(), node), vertices_.end()); } // root nodes are all vertices with 0 in-degrees std::vector ihipGraph::GetRootNodes() const { std::vector roots; for (auto entry : vertices_) { - if (nodeInDegree_.at(entry) == 0) { + if (entry->GetInDegree() == 0) { roots.push_back(entry); ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] root node: %s(%p)\n", GetGraphNodeTypeString(entry->GetType()), entry); @@ -101,7 +76,7 @@ std::vector ihipGraph::GetRootNodes() const { std::vector ihipGraph::GetLeafNodes() const { std::vector leafNodes; for (auto entry : vertices_) { - if (nodeOutDegree_.at(entry) == 0) { + if (entry->GetOutDegree() == 0) { leafNodes.push_back(entry); } } @@ -111,7 +86,7 @@ std::vector ihipGraph::GetLeafNodes() const { size_t ihipGraph::GetLeafNodeCount() const { int numLeafNodes = 0; for (auto entry : vertices_) { - if (nodeOutDegree_.at(entry) == 0) { + if (entry->GetOutDegree() == 0) { numLeafNodes++; } } @@ -120,23 +95,23 @@ size_t ihipGraph::GetLeafNodeCount() const { std::vector> ihipGraph::GetEdges() const { std::vector> edges; - for (const auto& i : edges_) { - for (const auto& j : i.second) { - edges.push_back(std::make_pair(i.first, j)); + for (const auto& i : vertices_) { + for (const auto& j : i->GetEdges()) { + edges.push_back(std::make_pair(i, j)); } } return edges; } void ihipGraph::GetRunListUtil(Node v, std::unordered_map& visited, - std::vector& singleList, - std::vector>& parallelLists, - std::unordered_map>& dependencies) { + std::vector& singleList, + std::vector>& parallelLists, + std::unordered_map>& dependencies) { // Mark the current node as visited. visited[v] = true; singleList.push_back(v); // Recurse for all the vertices adjacent to this vertex - for (auto& adjNode : edges_[v]) { + for (auto& adjNode : v->GetEdges()) { if (!visited[adjNode]) { // For the parallel list nodes add parent as the dependency if (singleList.empty()) { @@ -175,7 +150,7 @@ void ihipGraph::GetRunListUtil(Node v, std::unordered_map& visited, // The function to do Topological Sort. // It uses recursive GetRunListUtil() void ihipGraph::GetRunList(std::vector>& parallelLists, - std::unordered_map>& dependencies) { + std::unordered_map>& dependencies) { std::vector singleList; // Mark all the vertices as not visited @@ -184,6 +159,8 @@ void ihipGraph::GetRunList(std::vector>& parallelLists, // Call the recursive helper function for all vertices one by one for (auto node : vertices_) { + // If the node has embedded child graph + node->GetRunList(parallelLists, dependencies); if (visited[node] == false) { GetRunListUtil(node, visited, singleList, parallelLists, dependencies); } @@ -196,7 +173,7 @@ void ihipGraph::GetRunList(std::vector>& parallelLists, } } -hipError_t ihipGraph::LevelOrder(std::vector& levelOrder) { +void ihipGraph::LevelOrder(std::vector& levelOrder) { std::vector roots = GetRootNodes(); std::unordered_map visited; std::queue q; @@ -209,7 +186,7 @@ hipError_t ihipGraph::LevelOrder(std::vector& levelOrder) { Node& node = q.front(); q.pop(); levelOrder.push_back(node); - for (const auto& i : edges_[node]) { + for (const auto& i : node->GetEdges()) { if (visited.find(i) == visited.end() && i->GetLevel() == (node->GetLevel() + 1)) { q.push(i); ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] %s(%p) level:%d \n", @@ -218,7 +195,36 @@ hipError_t ihipGraph::LevelOrder(std::vector& levelOrder) { } } } - return hipSuccess; +} + +ihipGraph* ihipGraph::clone() const { + ihipGraph* newGraph = new ihipGraph(); + std::unordered_map clonedNodes; + for (auto entry : vertices_) { + hipGraphNode* node = entry->clone(); + newGraph->vertices_.push_back(node); + clonedNodes[entry] = node; + } + std::vector dependancies; + std::vector clonedEdges; + std::vector clonedDependencies; + for (auto node : vertices_) { + const std::vector& edges = node->GetEdges(); + clonedEdges.clear(); + for (auto edge : edges) { + clonedEdges.push_back(clonedNodes[edge]); + } + clonedNodes[node]->SetEdges(clonedEdges); + } + for (auto node : vertices_) { + const std::vector& dependencies = node->GetDependencies(); + clonedDependencies.clear(); + for (auto dep : dependencies) { + clonedDependencies.push_back(clonedNodes[dep]); + } + clonedNodes[node]->SetDependencies(clonedDependencies); + } + return newGraph; } hipError_t hipGraphExec::CreateQueues() { @@ -253,25 +259,31 @@ hipError_t hipGraphExec::FillCommands() { } i++; } + + 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& node : levelOrder_) { - auto nodeWaitList = nodeWaitLists_.find(node); - if (nodeWaitList != nodeWaitLists_.end()) { - amd::Command::EventWaitList waitList; - for (auto depNode : nodeWaitList->second) { - for (auto command : depNode->GetCommands()) { - waitList.push_back(command); - } - } - for (auto command : nodeWaitList->first->GetCommands()) { - command->updateEventWaitList(waitList); + for (auto entry : nodeWaitLists_) { + amd::Command::EventWaitList waitList; + for (auto depNode : entry.second) { + for (auto command : depNode->GetCommands()) { + waitList.push_back(command); } } + for (auto command : entry.first->GetCommands()) { + command->updateEventWaitList(waitList); + } } return status; } - hipError_t hipGraphExec::Init() { hipError_t status; status = CreateQueues(); @@ -297,9 +309,7 @@ void hipGraphExec::ResetGraph(cl_event event, cl_int command_exec_status, void* hipGraphExec::activeGraphExec_[reinterpret_cast(user_data)]; if (graphExec != nullptr) { for (auto& node : graphExec->levelOrder_) { - for (auto& command : node->GetCommands()) { - command->resetStatus(CL_INT_MAX); - } + node->ResetStatus(); } graphExec->rootCommand_->resetStatus(CL_INT_MAX); graphExec->bExecPending_.store(false); @@ -308,7 +318,7 @@ void hipGraphExec::ResetGraph(cl_event event, cl_int command_exec_status, void* } } -hipError_t hipGraphExec::UpdateGraphToWaitOnRoot() { +void hipGraphExec::UpdateGraphToWaitOnRoot() { for (auto& singleList : parallelLists_) { amd::Command::EventWaitList waitList; waitList.push_back(rootCommand_); @@ -319,7 +329,6 @@ hipError_t hipGraphExec::UpdateGraphToWaitOnRoot() { } } } - return hipSuccess; } hipError_t hipGraphExec::Run(hipStream_t stream) { @@ -342,9 +351,7 @@ hipError_t hipGraphExec::Run(hipStream_t stream) { } rootCommand_->enqueue(); for (auto& node : levelOrder_) { - for (auto& command : node->GetCommands()) { - command->enqueue(); - } + node->EnqueueCommands(); } amd::Command* command = new amd::Marker(*queue, false, graphLastCmdWaitList_); @@ -362,7 +369,7 @@ hipError_t hipGraphExec::Run(hipStream_t stream) { command->enqueue(); // Add the new barrier to stall the stream, until the callback is done - amd::Command::EventWaitList eventWaitList;; + amd::Command::EventWaitList eventWaitList; eventWaitList.push_back(command); amd::Command* block_command = new amd::Marker(*queue, !kMarkerDisableFlush, eventWaitList); if (block_command == nullptr) { diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 312aca3c14..cfe194aaf1 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -33,45 +33,161 @@ typedef hipGraphNode* Node; hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams); - struct hipGraphNode { protected: uint32_t level_; + unsigned int id_; hipGraphNodeType type_; std::vector commands_; + std::vector edges_; + std::vector dependencies_; bool visited_; + // count of in coming edges + size_t inDegree_; + // count of outgoing edges + size_t outDegree_; + static int nextID; + struct ihipGraph* parentGraph_; public: hipGraphNode(hipGraphNodeType type) { type_ = type; level_ = 0; visited_ = false; + inDegree_ = 0; + outDegree_ = 0; + id_ = nextID++; + parentGraph_ = nullptr; + } + /// Copy Constructor + hipGraphNode(const hipGraphNode& node) { + level_ = node.level_; + type_ = node.type_; + inDegree_ = node.inDegree_; + outDegree_ = node.outDegree_; + visited_ = false; + id_ = node.id_; + parentGraph_ = nullptr; } virtual ~hipGraphNode() { - for (auto command : commands_) { - delete command; + for (auto node : edges_) { + node->RemoveDependency(this); + } + for (auto node : dependencies_) { + node->RemoveEdge(this); } } + /// Create amd::command for the graph node virtual hipError_t CreateCommand(amd::HostQueue* queue) { return hipSuccess; } + /// Method to release amd::command part of node + virtual void ReleaseCommand() { + for (auto command : commands_) { + command->release(); + } + commands_.clear(); + } + /// Return node unique ID + int GetID() const { return id_; } + /// Returns command for graph node std::vector& GetCommands() { return commands_; } - hipGraphNodeType GetType() { return type_; } - uint32_t GetLevel() { return level_; } + /// Returns graph node type + hipGraphNodeType GetType() const { return type_; } + /// Returns graph node in coming edges + uint32_t GetLevel() const { return level_; } + /// Set graph node level void SetLevel(uint32_t level) { level_ = level; } + /// Clone graph node + virtual hipGraphNode* clone() const = 0; + /// Returns graph node indegree + size_t GetInDegree() const { return inDegree_; } + /// Updates indegree of the node + void SetInDegree(size_t inDegree) { inDegree_ = inDegree; } + /// Returns graph node outdegree + size_t GetOutDegree() const { return outDegree_; } + /// Updates outdegree of the node + void SetOutDegree(size_t outDegree) { outDegree_ = outDegree; } + /// Returns graph node dependencies + const std::vector& GetDependencies() const { return dependencies_; } + /// Update graph node dependecies + void SetDependencies(std::vector& dependencies) { + for (auto entry : dependencies) { + dependencies_.push_back(entry); + } + } + /// Add graph node dependency + void AddDependency(const Node& node) { dependencies_.push_back(node); } + /// Remove graph node dependency + void RemoveDependency(const Node& node) { + dependencies_.erase(std::remove(dependencies_.begin(), dependencies_.end(), node), + dependencies_.end()); + } + /// Return graph node children + const std::vector& GetEdges() const { return edges_; } + /// Updates graph node children + void SetEdges(std::vector& edges) { + for (auto entry : edges) { + edges_.push_back(entry); + } + } + /// Add edge, update parent node outdegree, child node indegree, level and dependency + void AddEdge(const Node& childNode) { + edges_.push_back(childNode); + outDegree_++; + childNode->SetInDegree(childNode->GetInDegree() + 1); + childNode->SetLevel(std::max(childNode->GetLevel(), GetLevel() + 1)); + childNode->AddDependency(this); + } + /// Remove edge, update parent node outdegree, child node indegree, level and dependency + void RemoveEdge(const Node& childNode) { + edges_.erase(std::remove(edges_.begin(), edges_.end(), childNode), edges_.end()); + outDegree_--; + childNode->SetInDegree(childNode->GetInDegree() - 1); + const std::vector& dependencies = childNode->GetDependencies(); + int32_t level = 0; + int32_t parentLevel = 0; + for (auto parent : dependencies) { + parentLevel = parent->GetLevel(); + level = std::max(level, (parentLevel + 1)); + } + childNode->SetLevel(level); + childNode->RemoveDependency(this); + } + /// Get Runlist of the nodes embedded as part of the graphnode(e.g. ChildGraph) + virtual void GetRunList(std::vector>& parallelList, + std::unordered_map>& dependencies) {} + /// 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() {} + /// Enqueue commands part of the node + virtual void EnqueueCommands() { + for (auto& command : commands_) { + command->enqueue(); + } + } + /// Reset commands part of the node + virtual void ResetStatus() { + for (auto& command : commands_) { + command->resetStatus(CL_INT_MAX); + } + } + ihipGraph* GetParentGraph() { return parentGraph_; } + void SetParentGraph(ihipGraph* graph) { parentGraph_ = graph; } }; struct ihipGraph { - std::unordered_map nodeInDegree_; // count of in coming edges for every vertex - std::unordered_map nodeOutDegree_; // count of outgoing edges for every vertex std::vector vertices_; - std::unordered_map> edges_; public: ihipGraph() {} - ~ihipGraph(){}; + ~ihipGraph() { + for (auto node : vertices_) { + delete node; + } + } /// add node to the graph - hipError_t AddNode(const Node& node); - /// add edge to the graph - hipError_t AddEdge(const Node& parentNode, const Node& childNode); + void AddNode(const Node& node); + void RemoveNode(const Node& node); /// Returns root nodes, all vertices with 0 in-degrees std::vector GetRootNodes() const; /// Returns leaf nodes, all vertices with 0 out-degrees @@ -81,7 +197,7 @@ struct ihipGraph { /// Returns total numbers of nodes in the graph size_t GetNodeCount() const { return vertices_.size(); } /// returns all the nodes in the graph - std::vector GetNodes() const { return vertices_; } + const std::vector& GetNodes() const { return vertices_; } /// returns all the edges in the graph std::vector> GetEdges() const; void GetRunListUtil(Node v, std::unordered_map& visited, @@ -89,7 +205,101 @@ struct ihipGraph { std::unordered_map>& dependencies); void GetRunList(std::vector>& parallelList, std::unordered_map>& dependencies); - hipError_t LevelOrder(std::vector& levelOrder); + void LevelOrder(std::vector& levelOrder); + ihipGraph* clone() const; +}; + +struct hipChildGraphNode : public hipGraphNode { + struct ihipGraph* childGraph_; + std::vector childGraphlevelOrder_; + + public: + hipChildGraphNode(ihipGraph* g) : hipGraphNode(hipGraphNodeTypeGraph) { + // ToDo: clone the child graph + childGraph_ = g->clone(); + } + + ~hipChildGraphNode() { delete childGraph_; } + + hipChildGraphNode(const hipChildGraphNode& rhs) : hipGraphNode(rhs) { + childGraph_ = rhs.childGraph_->clone(); + } + + hipGraphNode* clone() const { + return new hipChildGraphNode(static_cast(*this)); + } + + 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); + } + for (auto& node : childGraphlevelOrder_) { + node->ResetStatus(); + } + } + + void GetRunList(std::vector>& parallelList, + std::unordered_map>& dependencies) { + childGraph_->GetRunList(parallelList, dependencies); + } + + void LevelOrder(std::vector& levelOrder) { childGraph_->LevelOrder(levelOrder); } + + void EnqueueCommands() { + if (commands_.size() == 2) { + // enqueue child graph start command + commands_[0]->enqueue(); + // enqueue nodes in child graph in level order + for (auto& node : childGraphlevelOrder_) { + node->EnqueueCommands(); + } + // enqueue child graph end command + commands_[1]->enqueue(); + } else { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, + "[hipGraph] childgraph node commands are not created!\n"); + } + } }; class hipGraphKernelNode : public hipGraphNode { @@ -103,7 +313,13 @@ class hipGraphKernelNode : public hipGraphNode { func_ = func; } ~hipGraphKernelNode() { delete pKernelParams_; } - + hipGraphKernelNode(const hipGraphKernelNode& rhs) : hipGraphNode(rhs) { + pKernelParams_ = new hipKernelNodeParams(*rhs.pKernelParams_); + func_ = rhs.func_; + } + hipGraphNode* clone() const { + return new hipGraphKernelNode(static_cast(*this)); + } hipError_t CreateCommand(amd::HostQueue* queue) { commands_.reserve(1); amd::Command* command; @@ -153,6 +369,14 @@ class hipGraphMemcpyNode : public hipGraphNode { } ~hipGraphMemcpyNode() { delete pCopyParams_; } + hipGraphMemcpyNode(const hipGraphMemcpyNode& rhs) : hipGraphNode(rhs) { + pCopyParams_ = new hipMemcpy3DParms(*rhs.pCopyParams_); + } + + hipGraphNode* clone() const { + return new hipGraphMemcpyNode(static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue) { commands_.reserve(1); amd::Command* command; @@ -178,8 +402,13 @@ class hipGraphMemcpyNode1D : public hipGraphNode { public: hipGraphMemcpyNode1D(void* dst, const void* src, size_t count, hipMemcpyKind kind) : hipGraphNode(hipGraphNodeTypeMemcpy1D), dst_(dst), src_(src), count_(count), kind_(kind) {} + ~hipGraphMemcpyNode1D() {} + hipGraphNode* clone() const { + return new hipGraphMemcpyNode1D(static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue) { commands_.reserve(1); amd::Command* command = nullptr; @@ -214,6 +443,11 @@ template class hipGraphMemcpyNodeFromSymbol : public hipGraphNode { kind_(kind) {} ~hipGraphMemcpyNodeFromSymbol() {} + hipGraphNode* clone() const { + return new hipGraphMemcpyNodeFromSymbol( + static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue); void SetParams(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind) { @@ -243,6 +477,10 @@ template class hipGraphMemcpyNodeToSymbol : public hipGraphNode { kind_(kind) {} ~hipGraphMemcpyNodeToSymbol() {} + hipGraphNode* clone() const { + return new hipGraphMemcpyNodeToSymbol(static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue); void SetParams(const T& symbol, void* src, size_t count, size_t offset, hipMemcpyKind kind) { @@ -262,6 +500,14 @@ class hipGraphMemsetNode : public hipGraphNode { pMemsetParams_ = new hipMemsetParams(*pMemsetParams); } ~hipGraphMemsetNode() { delete pMemsetParams_; } + // Copy constructor + hipGraphMemsetNode(const hipGraphMemsetNode& memsetNode) : hipGraphNode(memsetNode) { + pMemsetParams_ = new hipMemsetParams(*memsetNode.pMemsetParams_); + } + + hipGraphNode* clone() const { + return new hipGraphMemsetNode(static_cast(*this)); + } hipError_t CreateCommand(amd::HostQueue* queue) { if (pMemsetParams_->height == 1) { @@ -294,6 +540,10 @@ class hipGraphEventRecordNode : public hipGraphNode { : hipGraphNode(hipGraphNodeTypeEventRecord), event_(event) {} ~hipGraphEventRecordNode() {} + hipGraphNode* clone() const { + return new hipGraphEventRecordNode(static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue); void GetParams(hipEvent_t* event) { *event = event_; } @@ -308,6 +558,10 @@ class hipGraphEventWaitNode : public hipGraphNode { : hipGraphNode(hipGraphNodeTypeWaitEvent), event_(event) {} ~hipGraphEventWaitNode() {} + hipGraphNode* clone() const { + return new hipGraphEventWaitNode(static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue); void GetParams(hipEvent_t* event) { *event = event_; } @@ -323,6 +577,14 @@ class hipGraphHostNode : public hipGraphNode { } ~hipGraphHostNode() { delete pNodeParams_; } + hipGraphHostNode(const hipGraphHostNode& hostNode) : hipGraphNode(hostNode) { + pNodeParams_ = new hipHostNodeParams(*hostNode.pNodeParams_); + } + + hipGraphNode* clone() const { + return new hipGraphHostNode(static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue); void GetParams(hipHostNodeParams* params) { @@ -338,6 +600,10 @@ class hipGraphEmptyNode : public hipGraphNode { hipGraphEmptyNode() : hipGraphNode(hipGraphNodeTypeEmpty) {} ~hipGraphEmptyNode() {} + hipGraphNode* clone() const { + return new hipGraphEmptyNode(static_cast(*this)); + } + hipError_t CreateCommand(amd::HostQueue* queue) { amd::Command::EventWaitList waitList; commands_.reserve(1); @@ -349,6 +615,7 @@ class hipGraphEmptyNode : public hipGraphNode { 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_; @@ -374,14 +641,14 @@ struct hipGraphExec { queue->release(); } for (auto node : levelOrder_) { - delete node; + node->ReleaseCommand(); } } hipError_t CreateQueues(); hipError_t FillCommands(); hipError_t Init(); - hipError_t UpdateGraphToWaitOnRoot(); + void UpdateGraphToWaitOnRoot(); hipError_t Run(hipStream_t stream); static void ResetGraph(cl_event event, cl_int command_exec_status, void* user_data); }; diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index 437a0c9d05..878714b7ca 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -177,11 +177,13 @@ namespace hip { hipStream_t parentStream_; /// Last graph node captured in the stream std::vector lastCapturedNodes_; + /// dependencies removed via API hipStreamUpdateCaptureDependencies + std::vector removedDependencies_; /// Derived streams/Paralell branches from the origin stream std::vector parallelCaptureStreams_; /// Capture events std::vector captureEvents_; - + unsigned long long captureID_; public: Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false, const std::vector& cuMask = {}, @@ -231,8 +233,19 @@ namespace hip { lastCapturedNodes_.clear(); lastCapturedNodes_.push_back(graphNode); } + /// returns updated dependencies removed + const std::vector& GetRemovedDependencies() { + return removedDependencies_; + } /// Append captured node via the wait event cross stream - void AddCrossCapturedNode(std::vector graphNodes) { + void AddCrossCapturedNode(std::vector graphNodes, bool replace = false) { + // replace dependencies as per flag hipStreamSetCaptureDependencies + if (replace == true) { + for (auto node : lastCapturedNodes_) { + removedDependencies_.push_back(node); + } + lastCapturedNodes_.clear(); + } for (auto node : graphNodes) { lastCapturedNodes_.push_back(node); } @@ -241,6 +254,8 @@ namespace hip { void SetCaptureGraph(hipGraph_t pGraph) { pCaptureGraph_ = pGraph; captureStatus_ = hipStreamCaptureStatusActive; + // ID is generated in Begin Capture i.e.. when capture status is active + captureID_ = GenerateCaptureID(); } /// reset capture parameters hipError_t EndCapture(); @@ -252,6 +267,13 @@ namespace hip { void SetParentStream(hipStream_t parentStream) { parentStream_ = parentStream; } /// Get parent stream hipStream_t GetParentStream() { return parentStream_; } + /// Generate ID for stream capture unique over the lifetime of the process + static int GenerateCaptureID() { + static std::atomic uid(0); + return ++uid; + } + /// Get Capture ID + int GetCaptureID() { return captureID_; } }; /// HIP Device class diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 280f132e0a..4b4d62f779 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -41,7 +41,8 @@ Stream::Stream(hip::Device* dev, Priority p, unsigned int f, bool null_stream, null_(null_stream), cuMask_(cuMask), captureStatus_(captureStatus), - originStream_(false) {} + originStream_(false), + captureID_(0) {} // ================================================================================================ Stream::~Stream() {