From d8022cd0811d206a5721799da82fe709843bb125 Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Tue, 25 Jul 2023 12:54:23 +0000 Subject: [PATCH] SWDEV-407568 - Move graph implementation to hip namespace Change-Id: I7023f202a7e3eb25b17db6d3e361205594ae81a5 [ROCm/clr commit: fd97dde1e62b50d2d9d56c2ab2194bf603359333] --- projects/clr/hipamd/src/hip_event.cpp | 2 +- projects/clr/hipamd/src/hip_event.hpp | 6 +- projects/clr/hipamd/src/hip_graph.cpp | 701 +++++++++++------- .../clr/hipamd/src/hip_graph_internal.cpp | 66 +- .../clr/hipamd/src/hip_graph_internal.hpp | 367 ++++----- projects/clr/hipamd/src/hip_internal.hpp | 25 +- 6 files changed, 653 insertions(+), 514 deletions(-) diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 50e5c091b8..5c746f1dc0 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -399,7 +399,7 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream) { e->SetCaptureStream(stream); if ((s != nullptr) && (s->GetCaptureStatus() == hipStreamCaptureStatusActive)) { s->SetCaptureEvent(event); - std::vector lastCapturedNodes = s->GetLastCapturedNodes(); + std::vector lastCapturedNodes = s->GetLastCapturedNodes(); if (!lastCapturedNodes.empty()) { e->SetNodesPrevToRecorded(lastCapturedNodes); } diff --git a/projects/clr/hipamd/src/hip_event.hpp b/projects/clr/hipamd/src/hip_event.hpp index bb6f6a239c..df32e37262 100644 --- a/projects/clr/hipamd/src/hip_event.hpp +++ b/projects/clr/hipamd/src/hip_event.hpp @@ -94,7 +94,7 @@ class Event { /// capture stream where event is recorded hipStream_t captureStream_ = nullptr; /// Previous captured nodes before event record - std::vector nodesPrevToRecorded_; + std::vector nodesPrevToRecorded_; protected: bool CheckHwEvent(eventType type) { bool ready; @@ -153,9 +153,9 @@ class Event { /// Set capture stream where event is recorded void SetCaptureStream(hipStream_t stream) { captureStream_ = stream; } /// Returns previous captured nodes before event record - std::vector GetNodesPrevToRecorded() const { return nodesPrevToRecorded_; } + std::vector GetNodesPrevToRecorded() const { return nodesPrevToRecorded_; } /// Set last captured graph node before event record - void SetNodesPrevToRecorded(std::vector& graphNode) { + void SetNodesPrevToRecorded(std::vector& graphNode) { nodesPrevToRecorded_ = graphNode; } virtual hipError_t GetHandle(ihipIpcEventHandle_t* handle) { diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 62e5642448..8db887607b 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -31,13 +31,13 @@ amd::Monitor g_captureStreamsLock{"StreamCaptureGlobalList"}; amd::Monitor g_streamSetLock{"StreamCaptureset"}; std::unordered_set g_allCapturingStreams; -inline hipError_t ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, +inline hipError_t ihipGraphAddNode(hip::GraphNode* graphNode, hip::Graph* graph, + hip::GraphNode* const* pDependencies, size_t numDependencies, bool capture = true) { graph->AddNode(graphNode); - std::unordered_set DuplicateDep; + std::unordered_set DuplicateDep; for (size_t i = 0; i < numDependencies; i++) { - if ((!hipGraphNode::isNodeValid(pDependencies[i])) || + if ((!hip::GraphNode::isNodeValid(pDependencies[i])) || (graph != pDependencies[i]->GetParentGraph())) { return hipErrorInvalidValue; } @@ -61,8 +61,8 @@ inline hipError_t ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph, return hipSuccess; } -hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, +hipError_t ihipGraphAddKernelNode(hip::GraphNode** pGraphNode, hip::Graph* graph, + hip::GraphNode* const* pDependencies, size_t numDependencies, const hipKernelNodeParams* pNodeParams, bool capture = true) { if (pGraphNode == nullptr || graph == nullptr || (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr || @@ -70,7 +70,7 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, return hipErrorInvalidValue; } - if (!ihipGraph::isGraphValid(graph)) { + if (!hip::Graph::isGraphValid(graph)) { return hipErrorInvalidValue; } @@ -80,7 +80,7 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, return hipErrorInvalidValue; } - hipError_t status = hipGraphKernelNode::validateKernelParams(pNodeParams); + hipError_t status = hip::GraphKernelNode::validateKernelParams(pNodeParams); if (hipSuccess != status) { return status; } @@ -94,13 +94,13 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, return hipErrorInvalidConfiguration; } - *pGraphNode = new hipGraphKernelNode(pNodeParams); + *pGraphNode = new hip::GraphKernelNode(pNodeParams); status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture); return status; } -hipError_t ihipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, +hipError_t ihipGraphAddMemcpyNode(hip::GraphNode** pGraphNode, hip::Graph* graph, + hip::GraphNode* const* pDependencies, size_t numDependencies, const hipMemcpy3DParms* pCopyParams, bool capture = true) { if (pGraphNode == nullptr || graph == nullptr || (numDependencies > 0 && pDependencies == nullptr) || pCopyParams == nullptr) { @@ -110,30 +110,30 @@ hipError_t ihipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, if (status != hipSuccess) { return status; } - *pGraphNode = new hipGraphMemcpyNode(pCopyParams); + *pGraphNode = new hip::GraphMemcpyNode(pCopyParams); status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture); return status; } -hipError_t ihipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, +hipError_t ihipGraphAddMemcpyNode1D(hip::GraphNode** pGraphNode, hip::Graph* graph, + hip::GraphNode* const* pDependencies, size_t numDependencies, void* dst, const void* src, size_t count, hipMemcpyKind kind, bool capture = true) { if (pGraphNode == nullptr || graph == nullptr || (numDependencies > 0 && pDependencies == nullptr) || count ==0) { return hipErrorInvalidValue; } - hipError_t status = hipGraphMemcpyNode1D::ValidateParams(dst, src, count, kind); + hipError_t status = hip::GraphMemcpyNode1D::ValidateParams(dst, src, count, kind); if (status != hipSuccess) { return status; } - *pGraphNode = new hipGraphMemcpyNode1D(dst, src, count, kind); + *pGraphNode = new hip::GraphMemcpyNode1D(dst, src, count, kind); status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture); return status; } -hipError_t ihipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, +hipError_t ihipGraphAddMemsetNode(hip::GraphNode** pGraphNode, hip::Graph* graph, + hip::GraphNode* const* pDependencies, size_t numDependencies, const hipMemsetParams* pMemsetParams, bool capture = true) { if (pGraphNode == nullptr || graph == nullptr || pMemsetParams == nullptr || (numDependencies > 0 && pDependencies == nullptr) || pMemsetParams->height == 0) { @@ -167,7 +167,7 @@ hipError_t ihipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, if (status != hipSuccess) { return status; } - *pGraphNode = new hipGraphMemsetNode(pMemsetParams); + *pGraphNode = new hip::GraphMemsetNode(pMemsetParams); status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture); return status; } @@ -189,7 +189,7 @@ hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction nodeParams.kernelParams = args; nodeParams.sharedMemBytes = sharedMemBytes; - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipError_t status = ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), &nodeParams); @@ -211,10 +211,10 @@ hipError_t ihipExtLaunchKernel(hipStream_t stream, hipFunction_t f, uint32_t glo } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipError_t status; if (startEvent != nullptr) { - pGraphNode = new hipGraphEventRecordNode(startEvent); + pGraphNode = new hip::GraphEventRecordNode(startEvent); status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), capture); if (status != hipSuccess) { @@ -239,7 +239,7 @@ hipError_t ihipExtLaunchKernel(hipStream_t stream, hipFunction_t f, uint32_t glo } s->SetLastCapturedNode(pGraphNode); if (stopEvent != nullptr) { - pGraphNode = new hipGraphEventRecordNode(stopEvent); + pGraphNode = new hip::GraphEventRecordNode(stopEvent); status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size()); if (status != hipSuccess) { @@ -294,7 +294,7 @@ hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, u nodeParams.kernelParams = kernelParams; nodeParams.sharedMemBytes = sharedMemBytes; - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipError_t status = ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), &nodeParams); @@ -312,7 +312,7 @@ hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*& return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipError_t status = ihipGraphAddMemcpyNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), p); @@ -335,7 +335,7 @@ hipError_t capturehipMemcpy2DAsync(hipStream_t& stream, void*& dst, size_t& dpit return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); @@ -373,7 +373,7 @@ hipError_t capturehipMemcpy2DFromArrayAsync(hipStream_t& stream, void*& dst, siz return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); p.srcPos = {wOffsetSrc, hOffsetSrc, 0}; @@ -408,7 +408,7 @@ hipError_t capturehipMemcpyFromArrayAsync(hipStream_t& stream, void*& dst, hipAr return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); p.srcPos = {wOffsetSrc, hOffsetSrc, 0}; @@ -446,7 +446,7 @@ hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray*& dst, s return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); p.dstPos = {wOffset, hOffset, 0}; @@ -481,7 +481,7 @@ hipError_t capturehipMemcpyToArrayAsync(hipStream_t& stream, hipArray_t& dst, si return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); p.dstPos = {wOffset, hOffset, 0}; @@ -514,7 +514,7 @@ hipError_t capturehipMemcpyParam2DAsync(hipStream_t& stream, const hip_Memcpy2D* return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); p.srcArray = pCopy->srcArray; @@ -566,7 +566,7 @@ hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArr return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); p.srcArray = srcArray; @@ -594,7 +594,7 @@ hipError_t capturehipMemcpyHtoAAsync(hipStream_t& stream, hipArray*& dstArray, s return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipMemcpy3DParms p = {}; memset(&p, 0, sizeof(p)); p.dstArray = dstArray; @@ -617,14 +617,14 @@ hipError_t capturehipMemcpy(hipStream_t stream, void* dst, const void* src, size return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); - std::vector pDependencies = s->GetLastCapturedNodes(); + std::vector pDependencies = s->GetLastCapturedNodes(); size_t numDependencies = s->GetLastCapturedNodes().size(); - hipGraph_t graph = s->GetCaptureGraph(); + hip::Graph* graph = s->GetCaptureGraph(); hipError_t status = ihipMemcpy_validate(dst, src, sizeBytes, kind); if (status != hipSuccess) { return status; } - hipGraphNode_t node = new hipGraphMemcpyNode1D(dst, src, sizeBytes, kind); + hip::GraphNode* node = new hip::GraphMemcpyNode1D(dst, src, sizeBytes, kind); status = ihipGraphAddNode(node, graph, pDependencies.data(), numDependencies); if (status != hipSuccess) { return status; @@ -689,8 +689,8 @@ hipError_t capturehipMemcpyFromSymbolAsync(hipStream_t& stream, void*& dst, cons HIP_RETURN(status); } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode = - new hipGraphMemcpyNodeFromSymbol(dst, symbol, sizeBytes, offset, kind); + hip::GraphNode* pGraphNode = + new hip::GraphMemcpyNodeFromSymbol(dst, symbol, sizeBytes, offset, kind); status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size()); if (status != hipSuccess) { @@ -714,7 +714,7 @@ hipError_t capturehipMemcpyToSymbolAsync(hipStream_t& stream, const void*& symbo HIP_RETURN(status); } hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode = new hipGraphMemcpyNodeToSymbol(symbol, src, sizeBytes, offset, kind); + hip::GraphNode* pGraphNode = new hip::GraphMemcpyNodeToSymbol(symbol, src, sizeBytes, offset, kind); status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size()); if (status != hipSuccess) { @@ -739,7 +739,7 @@ hipError_t capturehipMemsetAsync(hipStream_t& stream, void*& dst, int& value, si memsetParams.height = 1; hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipError_t status = ihipGraphAddMemsetNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), &memsetParams); @@ -764,7 +764,7 @@ hipError_t capturehipMemset2DAsync(hipStream_t& stream, void*& dst, size_t& pitc memsetParams.height = height; memsetParams.pitch = pitch; hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode; + hip::GraphNode* pGraphNode; hipError_t status = ihipGraphAddMemsetNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), &memsetParams); @@ -798,7 +798,7 @@ hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*& hostParams.fn = fn; hostParams.userData = userData; hip::Stream* s = reinterpret_cast(stream); - hipGraphNode_t pGraphNode = new hipGraphHostNode(&hostParams); + hip::GraphNode* pGraphNode = new hip::GraphHostNode(&hostParams); hipError_t status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size()); @@ -833,7 +833,7 @@ hipError_t capturehipMallocAsync(hipStream_t stream, hipMemPool_t mem_pool, node_params.accessDescCount = descs.size(); node_params.bytesize = size; - auto mem_alloc_node = new hipGraphMemAllocNode(&node_params); + auto mem_alloc_node = new hip::GraphMemAllocNode(&node_params); auto status = ihipGraphAddNode(mem_alloc_node, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size()); if (status != hipSuccess) { @@ -849,7 +849,7 @@ hipError_t capturehipMallocAsync(hipStream_t stream, hipMemPool_t mem_pool, // ================================================================================================ hipError_t capturehipFreeAsync(hipStream_t stream, void* dev_ptr) { hip::Stream* s = reinterpret_cast(stream); - auto mem_free_node = new hipGraphMemFreeNode(dev_ptr); + auto mem_free_node = new hip::GraphMemFreeNode(dev_ptr); auto status = ihipGraphAddNode(mem_free_node, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size()); if (status != hipSuccess) { @@ -925,7 +925,7 @@ hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode return hipErrorIllegalState; } - s->SetCaptureGraph(new ihipGraph(s->GetDevice())); + s->SetCaptureGraph(new hip::Graph(s->GetDevice())); s->SetCaptureId(); s->SetCaptureMode(mode); s->SetOriginStream(); @@ -954,7 +954,7 @@ hipError_t hipStreamBeginCapture_spt(hipStream_t stream, hipStreamCaptureMode mo HIP_RETURN_DURATION(hipStreamBeginCapture_common(stream, mode)); } -hipError_t hipStreamEndCapture_common(hipStream_t stream, hipGraph_t* pGraph) { +hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) { if (pGraph == nullptr) { return hipErrorInvalidValue; } @@ -999,22 +999,22 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hipGraph_t* pGraph) { // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do // not result in hipErrorStreamCaptureUnjoined // add temporary node to check if all parallel streams have joined - hipGraphNode_t pGraphNode; - pGraphNode = new hipGraphEmptyNode(); + hip::GraphNode* pGraphNode; + pGraphNode = reinterpret_cast(new hip::GraphEmptyNode()); hipError_t status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size()); if (s->GetCaptureGraph()->GetLeafNodeCount() > 1) { - std::vector leafNodes = s->GetCaptureGraph()->GetLeafNodes(); - std::unordered_set nodes = s->GetCaptureGraph()->GetManualNodesDuringCapture(); + std::vector leafNodes = s->GetCaptureGraph()->GetLeafNodes(); + std::unordered_set nodes = s->GetCaptureGraph()->GetManualNodesDuringCapture(); for (auto node : nodes) { const auto& fnode = std::find(leafNodes.begin(), leafNodes.end(), node); if (fnode != leafNodes.end()) { leafNodes.erase(fnode); } } - const std::vector& removedDepNodes = s->GetRemovedDependencies(); + const std::vector& removedDepNodes = s->GetRemovedDependencies(); bool foundInRemovedDep = false; for (auto leafNode : leafNodes) { for (auto node : removedDepNodes) { @@ -1040,13 +1040,19 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hipGraph_t* pGraph) { hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { HIP_INIT_API(hipStreamEndCapture, stream, pGraph); - HIP_RETURN_DURATION(hipStreamEndCapture_common(stream, pGraph)); + if (pGraph == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + hip::Graph* graph; + hipError_t status = hipStreamEndCapture_common(stream, &graph); + *pGraph = reinterpret_cast(graph); + HIP_RETURN(status); } hipError_t hipStreamEndCapture_spt(hipStream_t stream, hipGraph_t* pGraph) { HIP_INIT_API(hipStreamEndCapture, stream, pGraph); PER_THREAD_DEFAULT_STREAM(stream); - HIP_RETURN_DURATION(hipStreamEndCapture_common(stream, pGraph)); + HIP_RETURN_DURATION(hipStreamEndCapture_common(stream, reinterpret_cast(pGraph))); } hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags) { @@ -1054,7 +1060,7 @@ hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags) { if ((pGraph == nullptr) || (flags != 0)) { HIP_RETURN(hipErrorInvalidValue); } - *pGraph = new ihipGraph(hip::getCurrentDevice()); + *pGraph = reinterpret_cast(new hip::Graph(hip::getCurrentDevice())); HIP_RETURN(hipSuccess); } @@ -1063,11 +1069,12 @@ hipError_t hipGraphDestroy(hipGraph_t graph) { if (graph == nullptr) { HIP_RETURN(hipErrorInvalidValue); } + hip::Graph* g = reinterpret_cast(graph); // if graph is not valid its destroyed already - if (!ihipGraph::isGraphValid(graph)) { + if (!hip::Graph::isGraphValid(g)) { HIP_RETURN(hipErrorIllegalState); } - delete graph; + delete g; HIP_RETURN(hipSuccess); } @@ -1080,9 +1087,12 @@ hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, (numDependencies > 0 && pDependencies == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } - - HIP_RETURN_DURATION(ihipGraphAddKernelNode(pGraphNode, graph, pDependencies, numDependencies, - pNodeParams, false)); + hip::GraphNode* node; + hipError_t status = ihipGraphAddKernelNode( + &node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), numDependencies, pNodeParams, false); + *pGraphNode = reinterpret_cast(node); + HIP_RETURN(status); } hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -1090,9 +1100,16 @@ hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const hipMemcpy3DParms* pCopyParams) { HIP_INIT_API(hipGraphAddMemcpyNode, pGraphNode, graph, pDependencies, numDependencies, pCopyParams); - - HIP_RETURN_DURATION(ihipGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, - pCopyParams, false)); + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + hip::GraphNode* node; + hipError_t status = ihipGraphAddMemcpyNode( + &node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), numDependencies, pCopyParams, false); + *pGraphNode = reinterpret_cast(node); + HIP_RETURN(status); } hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -1104,36 +1121,41 @@ hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, (numDependencies > 0 && pDependencies == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } - - HIP_RETURN_DURATION(ihipGraphAddMemcpyNode1D(pGraphNode, graph, pDependencies, numDependencies, - dst, src, count, kind, false)); + hip::GraphNode* node; + hipError_t status = ihipGraphAddMemcpyNode1D(&node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, dst, src, count, kind, false); + *pGraphNode = reinterpret_cast(node); + HIP_RETURN(status); } 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); - if (!hipGraphNode::isNodeValid(node) || dst == nullptr || src == nullptr || count == 0 || + hip::GraphNode* n = reinterpret_cast(node); + if (!hip::GraphNode::isNodeValid(n) || dst == nullptr || src == nullptr || count == 0 || src == dst) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(dst, src, count, kind)); + HIP_RETURN(reinterpret_cast(n)->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); - if (hGraphExec == nullptr || !hipGraphNode::isNodeValid(node) || dst == nullptr || + hip::GraphNode* n = reinterpret_cast(node); + if (hGraphExec == nullptr || !hip::GraphNode::isNodeValid(n) || dst == nullptr || src == nullptr || count == 0 || src == dst) { HIP_RETURN(hipErrorInvalidValue); } - - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast( + reinterpret_cast(hGraphExec)->GetClonedNode(n)); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(dst, src, count, kind)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(dst, src, count, kind)); } hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -1145,9 +1167,13 @@ hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, (numDependencies > 0 && pDependencies == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } - - HIP_RETURN_DURATION(ihipGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, - pMemsetParams, false)); + hip::GraphNode* node; + hipError_t status = + ihipGraphAddMemsetNode(&node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, pMemsetParams, false); + *pGraphNode = reinterpret_cast(node); + HIP_RETURN(status); } hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -1157,8 +1183,11 @@ hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, (numDependencies > 0 && pDependencies == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } - *pGraphNode = new hipGraphEmptyNode(); - hipError_t status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, false); + hip::GraphNode* node = new hip::GraphEmptyNode(); + hipError_t status = ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, false); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } @@ -1170,41 +1199,43 @@ hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t grap (numDependencies > 0 && pDependencies == nullptr) || childGraph == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *pGraphNode = new hipChildGraphNode(childGraph); - hipError_t status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, false); + hip::GraphNode* node = new hip::ChildGraphNode(reinterpret_cast(childGraph)); + hipError_t status = ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, false); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } -hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, - uint64_t flags = 0) { +hipError_t ihipGraphInstantiate(hip::GraphExec** pGraphExec, hip::Graph* graph, + uint64_t flags = 0) { if (pGraphExec == nullptr || graph == nullptr) { return hipErrorInvalidValue; } if (graph->IsGraphInstantiated() == true) { for (auto node : graph->GetNodes()) { - if ((node->GetType() == hipGraphNodeTypeMemAlloc) - || (node->GetType() == hipGraphNodeTypeMemFree)) { + if ((node->GetType() == hipGraphNodeTypeMemAlloc) || + (node->GetType() == hipGraphNodeTypeMemFree)) { return hipErrorNotSupported; } } } - std::unordered_map clonedNodes; - hipGraph_t clonedGraph = graph->clone(clonedNodes); + std::unordered_map clonedNodes; + hip::Graph* clonedGraph = graph->clone(clonedNodes); if (clonedGraph == nullptr) { return hipErrorInvalidValue; } - std::vector> parallelLists; - std::unordered_map> nodeWaitLists; - std::unordered_set graphExeUserObj; + std::vector> parallelLists; + std::unordered_map> nodeWaitLists; + std::unordered_set graphExeUserObj; clonedGraph->GetRunList(parallelLists, nodeWaitLists); - std::vector graphNodes; + std::vector graphNodes; if (false == clonedGraph->TopologicalOrder(graphNodes)) { return hipErrorInvalidValue; } clonedGraph->GetUserObjs(graphExeUserObj); - *pGraphExec = - new hipGraphExec(graphNodes, parallelLists, nodeWaitLists, clonedNodes, - graphExeUserObj, flags); + *pGraphExec = new hip::GraphExec(graphNodes, parallelLists, nodeWaitLists, clonedNodes, + graphExeUserObj, flags); if (*pGraphExec != nullptr) { graph->SetGraphInstantiated(true); return (*pGraphExec)->Init(); @@ -1216,7 +1247,13 @@ hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize) { HIP_INIT_API(hipGraphInstantiate, pGraphExec, graph); - HIP_RETURN_DURATION(ihipGraphInstantiate(pGraphExec, graph)); + if (pGraphExec == nullptr || graph == nullptr) { + return hipErrorInvalidValue; + } + hip::GraphExec* ge; + hipError_t status = ihipGraphInstantiate(&ge, reinterpret_cast(graph)); + *pGraphExec = reinterpret_cast(ge); + HIP_RETURN(status); } hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph, @@ -1231,7 +1268,10 @@ hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t g HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN_DURATION(ihipGraphInstantiate(pGraphExec, graph, flags)); + hip::GraphExec* ge; + hipError_t status = ihipGraphInstantiate(&ge, reinterpret_cast(graph), flags); + *pGraphExec = reinterpret_cast(ge); + HIP_RETURN(status); } hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) { @@ -1239,19 +1279,20 @@ hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) { if (pGraphExec == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - delete pGraphExec; + hip::GraphExec* ge = reinterpret_cast(pGraphExec); + delete ge; HIP_RETURN(hipSuccess); } -hipError_t ihipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) { +hipError_t ihipGraphLaunch(hip::GraphExec* graphExec, hipStream_t stream) { if (!hip::isValid(stream)) { return hipErrorContextIsDestroyed; } return graphExec->Run(stream); } -hipError_t hipGraphLaunch_common(hipGraphExec_t graphExec, hipStream_t stream) { - if (graphExec == nullptr || !hipGraphExec::isGraphExecValid(graphExec)) { +hipError_t hipGraphLaunch_common(hip::GraphExec* graphExec, hipStream_t stream) { + if (graphExec == nullptr || !hip::GraphExec::isGraphExecValid(graphExec)) { return hipErrorInvalidValue; } if (!hip::isValid(stream)) { @@ -1262,13 +1303,13 @@ hipError_t hipGraphLaunch_common(hipGraphExec_t graphExec, hipStream_t stream) { hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) { HIP_INIT_API(hipGraphLaunch, graphExec, stream); - HIP_RETURN_DURATION(hipGraphLaunch_common(graphExec, stream)); + HIP_RETURN_DURATION(hipGraphLaunch_common(reinterpret_cast(graphExec), stream)); } hipError_t hipGraphLaunch_spt(hipGraphExec_t graphExec, hipStream_t stream) { HIP_INIT_API(hipGraphLaunch, graphExec, stream); PER_THREAD_DEFAULT_STREAM(stream); - HIP_RETURN_DURATION(hipGraphLaunch_common(graphExec, stream)); + HIP_RETURN_DURATION(hipGraphLaunch_common(reinterpret_cast(graphExec), stream)); } hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes) { @@ -1276,8 +1317,8 @@ hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* num if (graph == nullptr || numNodes == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - std::vector graphNodes; - if (false == graph->TopologicalOrder(graphNodes)) { + std::vector graphNodes; + if (false == reinterpret_cast(graph)->TopologicalOrder(graphNodes)) { HIP_RETURN(hipErrorInvalidValue); } if (nodes == nullptr) { @@ -1285,11 +1326,11 @@ hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* num HIP_RETURN(hipSuccess); } else if (*numNodes <= graphNodes.size()) { for (int i = 0; i < *numNodes; i++) { - nodes[i] = graphNodes[i]; + nodes[i] = reinterpret_cast(graphNodes[i]); } } else { for (int i = 0; i < graphNodes.size(); i++) { - nodes[i] = graphNodes[i]; + nodes[i] = reinterpret_cast(graphNodes[i]); } for (int i = graphNodes.size(); i < *numNodes; i++) { nodes[i] = nullptr; @@ -1306,17 +1347,17 @@ hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, if (graph == nullptr || pNumRootNodes == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - const std::vector nodes = graph->GetRootNodes(); + const std::vector nodes = reinterpret_cast(graph)->GetRootNodes(); if (pRootNodes == nullptr) { *pNumRootNodes = nodes.size(); HIP_RETURN(hipSuccess); } else if (*pNumRootNodes <= nodes.size()) { for (int i = 0; i < *pNumRootNodes; i++) { - pRootNodes[i] = nodes[i]; + pRootNodes[i] = reinterpret_cast(nodes[i]); } } else { for (int i = 0; i < nodes.size(); i++) { - pRootNodes[i] = nodes[i]; + pRootNodes[i] = reinterpret_cast(nodes[i]); } for (int i = nodes.size(); i < *pNumRootNodes; i++) { pRootNodes[i] = nullptr; @@ -1328,28 +1369,29 @@ hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams) { HIP_INIT_API(hipGraphKernelNodeGetParams, node, pNodeParams); - if (!hipGraphNode::isNodeValid(node) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || + pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(pNodeParams); + reinterpret_cast(node)->GetParams(pNodeParams); HIP_RETURN(hipSuccess); } hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node, const hipKernelNodeParams* pNodeParams) { HIP_INIT_API(hipGraphKernelNodeSetParams, node, pNodeParams); - if (!hipGraphNode::isNodeValid(node) || pNodeParams == nullptr || pNodeParams->func == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || pNodeParams == nullptr || pNodeParams->func == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); + HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); } hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams) { HIP_INIT_API(hipGraphMemcpyNodeGetParams, node, pNodeParams); - if (!hipGraphNode::isNodeValid(node) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(pNodeParams); + reinterpret_cast(node)->GetParams(pNodeParams); HIP_RETURN(hipSuccess); } @@ -1363,7 +1405,7 @@ hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt attr != hipKernelNodeAttributeCooperative) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(hNode)->SetAttrParams(attr, value)); + HIP_RETURN(reinterpret_cast(hNode)->SetAttrParams(attr, value)); } hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr, @@ -1376,21 +1418,22 @@ hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt attr != hipKernelNodeAttributeCooperative) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(hNode)->GetAttrParams(attr, value)); + HIP_RETURN(reinterpret_cast(hNode)->GetAttrParams(attr, value)); } hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams) { HIP_INIT_API(hipGraphMemcpyNodeSetParams, node, pNodeParams); - if (!hipGraphNode::isNodeValid(node) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); + 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); - if (hGraphExec == nullptr || !hipGraphNode::isNodeValid(node)) { + hip::GraphNode* n = reinterpret_cast(node); + if (hGraphExec == nullptr || !hip::GraphNode::isNodeValid(reinterpret_cast(n))) { HIP_RETURN(hipErrorInvalidValue); } if (ihipMemcpy3D_validate(pNodeParams) != hipSuccess) { @@ -1401,49 +1444,51 @@ hipError_t hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo ((pNodeParams->dstArray == 0) && (pNodeParams->dstPtr.ptr == nullptr))) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams)); } hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams) { HIP_INIT_API(hipGraphMemsetNodeGetParams, node, pNodeParams); - if (!hipGraphNode::isNodeValid(node) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(pNodeParams); + reinterpret_cast(node)->GetParams(pNodeParams); HIP_RETURN(hipSuccess); } hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams) { HIP_INIT_API(hipGraphMemsetNodeSetParams, node, pNodeParams); - if (!hipGraphNode::isNodeValid(node) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } if (pNodeParams->height > 1 && pNodeParams->pitch < (pNodeParams->width * pNodeParams->elementSize)) { return hipErrorInvalidValue; } - HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); + HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); } hipError_t hipGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, const hipMemsetParams* pNodeParams) { HIP_INIT_API(hipGraphExecMemsetNodeSetParams, hGraphExec, node, pNodeParams); - if (hGraphExec == nullptr || !hipGraphNode::isNodeValid(node) || pNodeParams == nullptr || + hip::GraphNode* n = reinterpret_cast(node); + + if (hGraphExec == nullptr || !hip::GraphNode::isNodeValid(n) || pNodeParams == nullptr || pNodeParams->dst == nullptr) { HIP_RETURN(hipErrorInvalidValue); } if (ihipGraphMemsetParams_validate(pNodeParams) != hipSuccess) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams, true)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams, true)); } hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, @@ -1457,25 +1502,28 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, } else if (from == nullptr || to == nullptr) { HIP_RETURN(hipErrorInvalidValue); } + hip::GraphNode* const* fromNode = reinterpret_cast(from); + hip::GraphNode* const* toNode = reinterpret_cast(to); + hip::Graph* g = reinterpret_cast(graph); for (size_t i = 0; i < numDependencies; i++) { - // When the same node is specified for both from and to - if (from[i] == nullptr || to[i] == nullptr || from[i] == to[i] || - !hipGraphNode::isNodeValid(to[i]) || !hipGraphNode::isNodeValid(from[i]) || - // making sure the nodes blong to the graph - to[i]->GetParentGraph() != graph || from[i]->GetParentGraph() != graph) { + // When the same node is specified for both fromNode and toNode + if (fromNode[i] == nullptr || toNode[i] == nullptr || fromNode[i] == toNode[i] || + !hip::GraphNode::isNodeValid(toNode[i]) || !hip::GraphNode::isNodeValid(fromNode[i]) || + // making sure the nodes belong toNode the graph + toNode[i]->GetParentGraph() != g || fromNode[i]->GetParentGraph() != g) { HIP_RETURN(hipErrorInvalidValue); } } for (size_t i = 0; i < numDependencies; i++) { - // When the same edge added from->to return invalid value - const std::vector& edges = from[i]->GetEdges(); + // When the same edge added fromNode->toNode return invalid value + const std::vector& edges = fromNode[i]->GetEdges(); for (auto edge : edges) { - if (edge == to[i]) { + if (edge == toNode[i]) { HIP_RETURN(hipErrorInvalidValue); } } - from[i]->AddEdge(to[i]); + fromNode[i]->AddEdge(toNode[i]); } HIP_RETURN(hipSuccess); } @@ -1483,23 +1531,26 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, const hipKernelNodeParams* pNodeParams) { HIP_INIT_API(hipGraphExecKernelNodeSetParams, hGraphExec, node, pNodeParams); - if (hGraphExec == nullptr || !hipGraphNode::isNodeValid(node) || pNodeParams == nullptr || - pNodeParams->func == nullptr) { + hip::GraphNode* n = reinterpret_cast(node); + if (hGraphExec == nullptr || + !hip::GraphNode::isNodeValid(n) || + pNodeParams == nullptr || pNodeParams->func == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams)); } hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph) { HIP_INIT_API(hipGraphChildGraphNodeGetGraph, node, pGraph); - if (!hipGraphNode::isNodeValid(node) || pGraph == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || pGraph == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *pGraph = reinterpret_cast(node)->GetChildGraph(); + hip::Graph* g = reinterpret_cast(node)->GetChildGraph(); + *pGraph = reinterpret_cast(g); if (*pGraph == nullptr) { HIP_RETURN(hipErrorInvalidValue); } @@ -1509,39 +1560,41 @@ hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGrap hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, hipGraph_t childGraph) { HIP_INIT_API(hipGraphExecChildGraphNodeSetParams, hGraphExec, node, childGraph); - if (hGraphExec == nullptr || !hipGraphNode::isNodeValid(node) || childGraph == nullptr || - !ihipGraph::isGraphValid(childGraph)) { + hip::GraphNode* n = reinterpret_cast(node); + hip::Graph* cg = reinterpret_cast(childGraph); + if (hGraphExec == nullptr || !hip::GraphNode::isNodeValid(n) || childGraph == nullptr || + !hip::Graph::isGraphValid(cg)) { HIP_RETURN(hipErrorInvalidValue); } - if (childGraph == node->GetParentGraph()) { + if (cg == n->GetParentGraph()) { HIP_RETURN(hipErrorUnknown); } // Validate whether the topology of node and childGraph matches - std::vector childGraphNodes1; - node->TopologicalOrder(childGraphNodes1); + std::vector childGraphNodes1; + n->TopologicalOrder(childGraphNodes1); - std::vector childGraphNodes2; - childGraph->TopologicalOrder(childGraphNodes2); + std::vector childGraphNodes2; + cg->TopologicalOrder(childGraphNodes2); if (childGraphNodes1.size() != childGraphNodes2.size()) { HIP_RETURN(hipErrorUnknown); } // Validate if the node insertion order matches else { - for (std::vector::size_type i = 0; i != childGraphNodes1.size(); i++) { + for (std::vector::size_type i = 0; i != childGraphNodes1.size(); i++) { if (childGraphNodes1[i]->GetType() != childGraphNodes2[i]->GetType()) { HIP_RETURN(hipErrorUnknown); } } } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(childGraph)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(cg)); } hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream, @@ -1606,10 +1659,10 @@ hipError_t hipStreamGetCaptureInfo_v2_common(hipStream_t stream, *id_out = s->GetCaptureID(); } if (graph_out != nullptr) { - *graph_out = s->GetCaptureGraph(); + *graph_out = reinterpret_cast(s->GetCaptureGraph()); } if (dependencies_out != nullptr) { - *dependencies_out = s->GetLastCapturedNodes().data(); + *dependencies_out = reinterpret_cast(s->GetLastCapturedNodes().data()); } if (numDependencies_out != nullptr) { *numDependencies_out = s->GetLastCapturedNodes().size(); @@ -1647,23 +1700,24 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t HIP_RETURN(hipErrorContextIsDestroyed); } hip::Stream* s = reinterpret_cast(stream); + hip::GraphNode** deps = reinterpret_cast(dependencies); if (s->GetCaptureStatus() == hipStreamCaptureStatusNone) { HIP_RETURN(hipErrorIllegalState); } if ((s->GetCaptureGraph()->GetNodeCount() < numDependencies) || - (numDependencies > 0 && dependencies == nullptr) || + (numDependencies > 0 && deps == nullptr) || (flags != 0 && flags != hipStreamAddCaptureDependencies && flags != hipStreamSetCaptureDependencies)) { HIP_RETURN(hipErrorInvalidValue); } - std::vector depNodes; - const std::vector& graphNodes = s->GetCaptureGraph()->GetNodes(); + std::vector depNodes; + const std::vector& graphNodes = s->GetCaptureGraph()->GetNodes(); for (int i = 0; i < numDependencies; i++) { - if ((dependencies[i] == nullptr) || - std::find(std::begin(graphNodes), std::end(graphNodes), dependencies[i]) == std::end(graphNodes)) { + if ((deps[i] == nullptr) || + std::find(std::begin(graphNodes), std::end(graphNodes), deps[i]) == std::end(graphNodes)) { HIP_RETURN(hipErrorInvalidValue); } - depNodes.push_back(dependencies[i]); + depNodes.push_back(deps[i]); } if (flags == hipStreamAddCaptureDependencies) { s->AddCrossCapturedNode(depNodes); @@ -1680,9 +1734,12 @@ hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* fr if (graph == nullptr || (numDependencies > 0 && (from == nullptr || to == nullptr))) { HIP_RETURN(hipErrorInvalidValue); } + hip::GraphNode* const* fromNode = reinterpret_cast(from); + hip::GraphNode* const* toNode = reinterpret_cast(to); + hip::Graph* g = reinterpret_cast(graph); for (size_t i = 0; i < numDependencies; i++) { - if (to[i]->GetParentGraph() != graph || from[i]->GetParentGraph() != graph || - from[i]->RemoveUpdateEdge(to[i]) == false) { + if (toNode[i]->GetParentGraph() != g || fromNode[i]->GetParentGraph() != g || + fromNode[i]->RemoveUpdateEdge(toNode[i]) == false) { HIP_RETURN(hipErrorInvalidValue); } } @@ -1696,20 +1753,20 @@ hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode (to == nullptr && from != nullptr)) { HIP_RETURN(hipErrorInvalidValue); } - const std::vector> edges = graph->GetEdges(); + const std::vector> edges = reinterpret_cast(graph)->GetEdges(); // returns only the number of edges in numEdges when from and to are null if (from == nullptr && to == nullptr) { *numEdges = edges.size(); HIP_RETURN(hipSuccess); } else if (*numEdges <= edges.size()) { for (int i = 0; i < *numEdges; i++) { - from[i] = edges[i].first; - to[i] = edges[i].second; + from[i] = reinterpret_cast(edges[i].first); + to[i] = reinterpret_cast(edges[i].second); } } else { for (int i = 0; i < edges.size(); i++) { - from[i] = edges[i].first; - to[i] = edges[i].second; + from[i] = reinterpret_cast(edges[i].first); + to[i] = reinterpret_cast(edges[i].second); } // If numEdges > actual number of edges, the remaining entries in from and to will be set to // NULL @@ -1726,20 +1783,21 @@ hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDependencies, size_t* pNumDependencies) { HIP_INIT_API(hipGraphNodeGetDependencies, node, pDependencies, pNumDependencies); - if (!hipGraphNode::isNodeValid(node) || pNumDependencies == nullptr) { + hip::GraphNode* n = reinterpret_cast(node); + if (!hip::GraphNode::isNodeValid(n) || pNumDependencies == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - const std::vector& dependencies = node->GetDependencies(); + const std::vector& dependencies = n->GetDependencies(); if (pDependencies == NULL) { *pNumDependencies = dependencies.size(); HIP_RETURN(hipSuccess); } else if (*pNumDependencies <= dependencies.size()) { for (int i = 0; i < *pNumDependencies; i++) { - pDependencies[i] = dependencies[i]; + pDependencies[i] = reinterpret_cast(dependencies[i]); } } else { for (int i = 0; i < dependencies.size(); i++) { - pDependencies[i] = dependencies[i]; + pDependencies[i] = reinterpret_cast(dependencies[i]); } // pNumDependencies > actual number of dependencies, the remaining entries in pDependencies will // be set to NULL @@ -1754,20 +1812,21 @@ hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDep hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pDependentNodes, size_t* pNumDependentNodes) { HIP_INIT_API(hipGraphNodeGetDependentNodes, node, pDependentNodes, pNumDependentNodes); - if (!hipGraphNode::isNodeValid(node) || pNumDependentNodes == nullptr) { + hip::GraphNode* n = reinterpret_cast(node); + if (!hip::GraphNode::isNodeValid(n) || pNumDependentNodes == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - const std::vector& dependents = node->GetEdges(); + const std::vector& dependents = n->GetEdges(); if (pDependentNodes == NULL) { *pNumDependentNodes = dependents.size(); HIP_RETURN(hipSuccess); } else if (*pNumDependentNodes <= dependents.size()) { for (int i = 0; i < *pNumDependentNodes; i++) { - pDependentNodes[i] = dependents[i]; + pDependentNodes[i] = reinterpret_cast(dependents[i]); } } else { for (int i = 0; i < dependents.size(); i++) { - pDependentNodes[i] = dependents[i]; + pDependentNodes[i] = reinterpret_cast(dependents[i]); } // pNumDependentNodes > actual number of dependents, the remaining entries in pDependentNodes // will be set to NULL @@ -1781,29 +1840,33 @@ hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pD hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType) { HIP_INIT_API(hipGraphNodeGetType, node, pType); - if (!hipGraphNode::isNodeValid(node) || pType == nullptr) { + hip::GraphNode* n = reinterpret_cast(node); + + if (!hip::GraphNode::isNodeValid(reinterpret_cast(n)) || pType == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *pType = node->GetType(); + *pType = n->GetType(); HIP_RETURN(hipSuccess); } hipError_t hipGraphDestroyNode(hipGraphNode_t node) { HIP_INIT_API(hipGraphDestroyNode, node); - if (!hipGraphNode::isNodeValid(node)) { + hip::GraphNode* n = reinterpret_cast(node); + + if (!hip::GraphNode::isNodeValid(reinterpret_cast(n))) { HIP_RETURN(hipErrorInvalidValue); } // First remove all the edges both incoming and outgoing from node. - for(auto& edge : node->GetEdges()) { - node->RemoveUpdateEdge(edge); + for (auto& edge : n->GetEdges()) { + n->RemoveUpdateEdge(edge); } - const std::vector& dependencies = node->GetDependencies(); - for(auto& parent: dependencies) { - parent->RemoveEdge(node); + const std::vector& dependencies = n->GetDependencies(); + for (auto& parent : dependencies) { + parent->RemoveEdge(n); parent->SetOutDegree(parent->GetOutDegree() - 1); } // Remove the node from graph. - node->GetParentGraph()->RemoveNode(node); + n->GetParentGraph()->RemoveNode(n); HIP_RETURN(hipSuccess); } @@ -1813,10 +1876,11 @@ hipError_t hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph) { if (originalGraph == nullptr || pGraphClone == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - if (!ihipGraph::isGraphValid(originalGraph)) { + hip::Graph* g = reinterpret_cast(originalGraph); + if (!hip::Graph::isGraphValid(g)) { HIP_RETURN(hipErrorInvalidValue); } - *pGraphClone = originalGraph->clone(); + *pGraphClone = reinterpret_cast(g->clone()); HIP_RETURN(hipSuccess); } @@ -1826,13 +1890,14 @@ hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t origina if (pNode == nullptr || originalNode == nullptr || clonedGraph == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - if (clonedGraph->getOriginalGraph() == nullptr) { + hip::Graph* iClonedGraph = reinterpret_cast(clonedGraph); + if (iClonedGraph->getOriginalGraph() == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - for (auto node : clonedGraph->GetNodes()) { - if (node->GetID() == originalNode->GetID()) { - *pNode = node; + for (auto node : iClonedGraph->GetNodes()) { + if (node->GetID() == reinterpret_cast(originalNode)->GetID()) { + *pNode = reinterpret_cast(node); HIP_RETURN(hipSuccess); } } @@ -1845,9 +1910,10 @@ hipError_t hipGraphAddMemcpyNodeFromSymbol(hipGraphNode_t* pGraphNode, hipGraph_ size_t count, size_t offset, hipMemcpyKind kind) { HIP_INIT_API(hipGraphAddMemcpyNodeFromSymbol, pGraphNode, graph, pDependencies, numDependencies, dst, symbol, count, offset, kind); + hip::Graph* g = reinterpret_cast(graph); if (graph == nullptr || pGraphNode == nullptr || count == 0 || (numDependencies > 0 && pDependencies == nullptr) || dst == nullptr || - !ihipGraph::isGraphValid(graph)) { + !hip::Graph::isGraphValid(g)) { HIP_RETURN(hipErrorInvalidValue); } @@ -1858,8 +1924,10 @@ hipError_t hipGraphAddMemcpyNodeFromSymbol(hipGraphNode_t* pGraphNode, hipGraph_ if (status != hipSuccess) { HIP_RETURN(status); } - *pGraphNode = new hipGraphMemcpyNodeFromSymbol(dst, symbol, count, offset, kind); - status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, false); + hip::GraphNode* node = new hip::GraphMemcpyNodeFromSymbol(dst, symbol, count, offset, kind); + status = ihipGraphAddNode( + node, g, reinterpret_cast(pDependencies), numDependencies, false); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } @@ -1869,12 +1937,13 @@ hipError_t hipGraphMemcpyNodeSetParamsFromSymbol(hipGraphNode_t node, void* dst, if (symbol == nullptr) { HIP_RETURN(hipErrorInvalidSymbol); } - if (!hipGraphNode::isNodeValid(node) || dst == nullptr || count == 0 || symbol == dst) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || dst == nullptr || + count == 0 || symbol == dst) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(dst, symbol, count, - offset, kind)); + HIP_RETURN(reinterpret_cast(node)->SetParams( + dst, symbol, count, offset, kind)); } hipError_t hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node, @@ -1885,16 +1954,18 @@ hipError_t hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec, if (symbol == nullptr) { HIP_RETURN(hipErrorInvalidSymbol); } - if (hGraphExec == nullptr || !hipGraphNode::isNodeValid(node) || dst == nullptr || count == 0 || symbol == dst) { + hip::GraphNode* n = reinterpret_cast(node); + if (hGraphExec == nullptr || !hip::GraphNode::isNodeValid(n) || dst == nullptr || count == 0 || + symbol == dst) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } constexpr bool kCheckDeviceIsSame = true; - HIP_RETURN(reinterpret_cast(clonedNode) + HIP_RETURN(reinterpret_cast(clonedNode) ->SetParams(dst, symbol, count, offset, kind, kCheckDeviceIsSame)); } @@ -1906,7 +1977,7 @@ hipError_t hipGraphAddMemcpyNodeToSymbol(hipGraphNode_t* pGraphNode, hipGraph_t HIP_INIT_API(hipGraphAddMemcpyNodeToSymbol, pGraphNode, graph, pDependencies, numDependencies, symbol, src, count, offset, kind); if (pGraphNode == nullptr || graph == nullptr || src == nullptr || count == 0 || - !ihipGraph::isGraphValid(graph) || (pDependencies == nullptr && numDependencies > 0)) { + !hip::Graph::isGraphValid(reinterpret_cast(graph)) || (pDependencies == nullptr && numDependencies > 0)) { HIP_RETURN(hipErrorInvalidValue); } size_t sym_size = 0; @@ -1915,11 +1986,14 @@ hipError_t hipGraphAddMemcpyNodeToSymbol(hipGraphNode_t* pGraphNode, hipGraph_t if (status != hipSuccess) { HIP_RETURN(status); } - *pGraphNode = new hipGraphMemcpyNodeToSymbol(symbol, src, count, offset, kind); - if (*pGraphNode == nullptr) { + hip::GraphNode* node = new hip::GraphMemcpyNodeToSymbol(symbol, src, count, offset, kind); + if (node == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, false); + status = ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, false); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } @@ -1930,12 +2004,13 @@ hipError_t hipGraphMemcpyNodeSetParamsToSymbol(hipGraphNode_t node, const void* if (symbol == nullptr) { HIP_RETURN(hipErrorInvalidSymbol); } - if (!hipGraphNode::isNodeValid(node) || src == nullptr || count == 0 || symbol == src) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || src == nullptr || + count == 0 || symbol == src) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(symbol, src, count, - offset, kind)); + HIP_RETURN(reinterpret_cast(node)->SetParams(symbol, src, count, + offset, kind)); } @@ -1948,16 +2023,17 @@ hipError_t hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hi if (symbol == nullptr) { HIP_RETURN(hipErrorInvalidSymbol); } - if (hGraphExec == nullptr || src == nullptr || !hipGraphNode::isNodeValid(node) || count == 0 || src == symbol) { + hip::GraphNode* n = reinterpret_cast(node); + if (hGraphExec == nullptr || src == nullptr || !hip::GraphNode::isNodeValid(n) || count == 0 || src == symbol) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } constexpr bool kCheckDeviceIsSame = true; - HIP_RETURN(reinterpret_cast(clonedNode) + HIP_RETURN(reinterpret_cast(clonedNode) ->SetParams(symbol, src, count, offset, kind, kCheckDeviceIsSame)); } @@ -1970,40 +2046,49 @@ hipError_t hipGraphAddEventRecordNode(hipGraphNode_t* pGraphNode, hipGraph_t gra (numDependencies > 0 && pDependencies == nullptr) || event == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *pGraphNode = new hipGraphEventRecordNode(event); - hipError_t status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, false); + hip::GraphNode* node = new hip::GraphEventRecordNode(event); + hipError_t status = ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, false); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } hipError_t hipGraphEventRecordNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out) { HIP_INIT_API(hipGraphEventRecordNodeGetEvent, node, event_out); - if (!hipGraphNode::isNodeValid(node) || event_out == nullptr || node->GetType() != hipGraphNodeTypeEventRecord) { + hip::GraphNode* n = reinterpret_cast(node); + if (!hip::GraphNode::isNodeValid(n) || event_out == nullptr || + n->GetType() != hipGraphNodeTypeEventRecord) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(event_out); + reinterpret_cast(n)->GetParams(event_out); HIP_RETURN(hipSuccess); } hipError_t hipGraphEventRecordNodeSetEvent(hipGraphNode_t node, hipEvent_t event) { HIP_INIT_API(hipGraphEventRecordNodeSetEvent, node, event); - if (!hipGraphNode::isNodeValid(node) || event == nullptr || node->GetType() != hipGraphNodeTypeEventRecord) { + hip::GraphNode* n = reinterpret_cast(node); + + if (!hip::GraphNode::isNodeValid(reinterpret_cast(n)) || event == nullptr || + n->GetType() != hipGraphNodeTypeEventRecord) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(event)); + HIP_RETURN(reinterpret_cast(n)->SetParams(event)); } hipError_t hipGraphExecEventRecordNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, hipEvent_t event) { HIP_INIT_API(hipGraphExecEventRecordNodeSetEvent, hGraphExec, hNode, event); - if (hGraphExec == nullptr || hNode == nullptr || event == nullptr || - hNode->GetType() != hipGraphNodeTypeEventRecord) { + hip::GraphNode* n = reinterpret_cast(hNode); + if (hGraphExec == nullptr || n == nullptr || event == nullptr || + n->GetType() != hipGraphNodeTypeEventRecord) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(event)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(event)); } hipError_t hipGraphAddEventWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -2014,40 +2099,51 @@ hipError_t hipGraphAddEventWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph (numDependencies > 0 && pDependencies == nullptr) || event == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *pGraphNode = new hipGraphEventWaitNode(event); - hipError_t status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, false); + hip::GraphNode* node = new hip::GraphEventWaitNode(event); + hipError_t status = ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, false); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } hipError_t hipGraphEventWaitNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out) { HIP_INIT_API(hipGraphEventWaitNodeGetEvent, node, event_out); - if (!hipGraphNode::isNodeValid(node) || event_out == nullptr || node->GetType() != hipGraphNodeTypeWaitEvent) { + hip::GraphNode* n = reinterpret_cast(node); + + if (!hip::GraphNode::isNodeValid(reinterpret_cast(n)) || event_out == nullptr || + n->GetType() != hipGraphNodeTypeWaitEvent) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(event_out); + reinterpret_cast(n)->GetParams(event_out); HIP_RETURN(hipSuccess); } hipError_t hipGraphEventWaitNodeSetEvent(hipGraphNode_t node, hipEvent_t event) { HIP_INIT_API(hipGraphEventWaitNodeSetEvent, node, event); - if (!hipGraphNode::isNodeValid(node) || event == nullptr || node->GetType() != hipGraphNodeTypeWaitEvent) { + hip::GraphNode* n = reinterpret_cast(node); + + if (!hip::GraphNode::isNodeValid(reinterpret_cast(n)) || event == nullptr || + n->GetType() != hipGraphNodeTypeWaitEvent) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(event)); + HIP_RETURN(reinterpret_cast(n)->SetParams(event)); } hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, hipEvent_t event) { HIP_INIT_API(hipGraphExecEventWaitNodeSetEvent, hGraphExec, hNode, event); + hip::GraphNode* n = reinterpret_cast(hNode); + if (hGraphExec == nullptr || hNode == nullptr || event == nullptr || - (hNode->GetType() != hipGraphNodeTypeWaitEvent)) { + (n->GetType() != hipGraphNodeTypeWaitEvent)) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(event)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(event)); } hipError_t hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -2059,41 +2155,45 @@ hipError_t hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, HIP_RETURN(hipErrorInvalidValue); } - *pGraphNode = new hipGraphHostNode(pNodeParams); - hipError_t status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, false); + hip::GraphNode* node = new hip::GraphHostNode(pNodeParams); + hipError_t status = ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, false); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } hipError_t hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNodeParams) { HIP_INIT_API(hipGraphHostNodeGetParams, node, pNodeParams); - if (!hipGraphNode::isNodeValid(node) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(pNodeParams); + reinterpret_cast(node)->GetParams(pNodeParams); HIP_RETURN(hipSuccess); } hipError_t hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParams* pNodeParams) { HIP_INIT_API(hipGraphHostNodeSetParams, node, pNodeParams); if (pNodeParams == nullptr || pNodeParams->fn == nullptr || - !hipGraphNode::isNodeValid(node)) { + !hip::GraphNode::isNodeValid(reinterpret_cast(node))) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); + HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); } hipError_t hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, const hipHostNodeParams* pNodeParams) { HIP_INIT_API(hipGraphExecHostNodeSetParams, hGraphExec, node, pNodeParams); + hip::GraphNode* n = reinterpret_cast(node); if (hGraphExec == nullptr || pNodeParams == nullptr || pNodeParams->fn == nullptr || - !hipGraphNode::isNodeValid(node)) { + !hip::GraphNode::isNodeValid(n)) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + hip::GraphNode* clonedNode = reinterpret_cast(hGraphExec)->GetClonedNode(n); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams)); } hipError_t hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph, @@ -2106,18 +2206,19 @@ hipError_t hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph, HIP_RETURN(hipErrorInvalidValue); } - std::vector newGraphNodes; - hGraph->TopologicalOrder(newGraphNodes); - std::vector& oldGraphExecNodes = hGraphExec->GetNodes(); + std::vector newGraphNodes; + reinterpret_cast(hGraph)->TopologicalOrder(newGraphNodes); + std::vector& oldGraphExecNodes = + reinterpret_cast(hGraphExec)->GetNodes(); if (newGraphNodes.size() != oldGraphExecNodes.size()) { *updateResult_out = hipGraphExecUpdateErrorTopologyChanged; HIP_RETURN(hipErrorGraphExecUpdateFailure); } - for (std::vector::size_type i = 0; i != newGraphNodes.size(); i++) { + for (std::vector::size_type i = 0; i != newGraphNodes.size(); i++) { if (newGraphNodes[i]->GetType() == oldGraphExecNodes[i]->GetType()) { hipError_t status = oldGraphExecNodes[i]->SetParams(newGraphNodes[i]); if (status != hipSuccess) { - *hErrorNode_out = newGraphNodes[i]; + *hErrorNode_out = reinterpret_cast(newGraphNodes[i]); if (status == hipErrorInvalidDeviceFunction) { *updateResult_out = hipGraphExecUpdateErrorUnsupportedFunctionChange; } else if (status == hipErrorInvalidValue || status == hipErrorInvalidDevicePointer) { @@ -2128,7 +2229,7 @@ hipError_t hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph, HIP_RETURN(hipErrorGraphExecUpdateFailure); } } else { - *hErrorNode_out = newGraphNodes[i]; + *hErrorNode_out = reinterpret_cast(newGraphNodes[i]); *updateResult_out = hipGraphExecUpdateErrorNodeTypeChanged; HIP_RETURN(hipErrorGraphExecUpdateFailure); } @@ -2139,43 +2240,50 @@ hipError_t hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph, // ================================================================================================ hipError_t hipGraphAddMemAllocNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, - hipMemAllocNodeParams* pNodeParams) { - HIP_INIT_API(hipGraphAddMemAllocNode, pGraphNode, graph, - pDependencies, numDependencies, pNodeParams); + const hipGraphNode_t* pDependencies, size_t numDependencies, + hipMemAllocNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphAddMemAllocNode, pGraphNode, graph, pDependencies, numDependencies, + pNodeParams); if (pGraphNode == nullptr || graph == nullptr || (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - if (pNodeParams->bytesize == 0 || pNodeParams->poolProps.allocType != hipMemAllocationTypePinned - || pNodeParams->poolProps.location.type != hipMemLocationTypeDevice) { + if (pNodeParams->bytesize == 0 || + pNodeParams->poolProps.allocType != hipMemAllocationTypePinned || + pNodeParams->poolProps.location.type != hipMemLocationTypeDevice) { pNodeParams->dptr = nullptr; HIP_RETURN(hipErrorInvalidValue); } if (pNodeParams->poolProps.location.type == hipMemLocationTypeDevice) { - if (pNodeParams->poolProps.location.id < 0 || pNodeParams->poolProps.location.id >= g_devices.size()) { + if (pNodeParams->poolProps.location.id < 0 || + pNodeParams->poolProps.location.id >= g_devices.size()) { HIP_RETURN(hipErrorInvalidValue); } } // Clear the pointer to allocated memory because it may contain stale/uninitialized data pNodeParams->dptr = nullptr; - auto mem_alloc_node = new hipGraphMemAllocNode(pNodeParams); - *pGraphNode = mem_alloc_node; - auto status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); + auto mem_alloc_node = new hip::GraphMemAllocNode(pNodeParams); + hip::GraphNode* node = mem_alloc_node; + auto status = + ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), numDependencies); // The address must be provided during the node creation time pNodeParams->dptr = (HIP_MEM_POOL_USE_VM) ? mem_alloc_node->ReserveAddress() : mem_alloc_node->Execute(); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } // ================================================================================================ hipError_t hipGraphMemAllocNodeGetParams(hipGraphNode_t node, hipMemAllocNodeParams* pNodeParams) { HIP_INIT_API(hipGraphMemAllocNodeGetParams, node, pNodeParams); - if (node == nullptr || pNodeParams == nullptr || !hipGraphNode::isNodeValid(node) - || node->GetType() != hipGraphNodeTypeMemAlloc) { + hip::GraphNode* n = reinterpret_cast(node); + if (node == nullptr || pNodeParams == nullptr || + !hip::GraphNode::isNodeValid(reinterpret_cast(n)) || + n->GetType() != hipGraphNodeTypeMemAlloc) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(pNodeParams); + reinterpret_cast(n)->GetParams(pNodeParams); HIP_RETURN(hipSuccess); } @@ -2204,20 +2312,24 @@ hipError_t hipGraphAddMemFreeNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, } } - auto mem_free_node = new hipGraphMemFreeNode(dev_ptr); - *pGraphNode = mem_free_node; - auto status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); + auto mem_free_node = new hip::GraphMemFreeNode(dev_ptr); + hip::GraphNode* node = mem_free_node; + auto status = + ihipGraphAddNode(node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), numDependencies); + *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } // ================================================================================================ hipError_t hipGraphMemFreeNodeGetParams(hipGraphNode_t node, void* dev_ptr) { HIP_INIT_API(hipGraphMemFreeNodeGetParams, node, dev_ptr); - if (node == nullptr || dev_ptr == nullptr || !hipGraphNode::isNodeValid(node) - || node->GetType() != hipGraphNodeTypeMemFree) { + hip::GraphNode* n = reinterpret_cast(node); + if (node == nullptr || dev_ptr == nullptr || !hip::GraphNode::isNodeValid(n) + || n->GetType() != hipGraphNodeTypeMemFree) { HIP_RETURN(hipErrorInvalidValue); } - reinterpret_cast(node)->GetParams(reinterpret_cast(dev_ptr)); + reinterpret_cast(n)->GetParams(reinterpret_cast(dev_ptr)); HIP_RETURN(hipSuccess); } @@ -2298,11 +2410,12 @@ hipError_t hipUserObjectCreate(hipUserObject_t* object_out, void* ptr, hipHostFn HIP_RETURN(hipErrorInvalidValue); } - *object_out = new hipUserObject(destroy, ptr, flags); + hip::UserObject *object = new hip::UserObject(destroy, ptr, flags); //! Creating object adds one reference. if (initialRefcount > 1) { - (*object_out)->increaseRefCount(static_cast(initialRefcount - 1)); + object->increaseRefCount(static_cast(initialRefcount - 1)); } + *object_out = reinterpret_cast(object); HIP_RETURN(hipSuccess); } @@ -2311,14 +2424,15 @@ hipError_t hipUserObjectRelease(hipUserObject_t object, unsigned int count) { if (object == nullptr || count == 0 || count > INT_MAX) { HIP_RETURN(hipErrorInvalidValue); } - if (object->referenceCount() < count || !hipUserObject::isUserObjvalid(object)) { + hip::UserObject* userObject = reinterpret_cast(object); + if (userObject->referenceCount() < count || !hip::UserObject::isUserObjvalid(userObject)) { HIP_RETURN(hipSuccess); } //! If all the counts are gone not longer need the obj in the list - if (object->referenceCount() == count) { - hipUserObject::removeUSerObj(object); + if (userObject->referenceCount() == count) { + hip::UserObject::removeUSerObj(userObject); } - object->decreaseRefCount(count); + userObject->decreaseRefCount(count); HIP_RETURN(hipSuccess); } @@ -2327,10 +2441,11 @@ hipError_t hipUserObjectRetain(hipUserObject_t object, unsigned int count) { if (object == nullptr || count == 0 || count > INT_MAX) { HIP_RETURN(hipErrorInvalidValue); } - if (!hipUserObject::isUserObjvalid(object)) { + hip::UserObject* userObject = reinterpret_cast(object); + if (!hip::UserObject::isUserObjvalid(userObject)) { HIP_RETURN(hipSuccess); } - object->increaseRefCount(count); + userObject->increaseRefCount(count); HIP_RETURN(hipSuccess); } @@ -2342,7 +2457,9 @@ hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, un (flags != 0 && flags != hipGraphUserObjectMove)) { HIP_RETURN(hipErrorInvalidValue); } - if (!hipUserObject::isUserObjvalid(object) && !graph->isUserObjGraphValid(object)) { + hip::UserObject* userObject = reinterpret_cast(object); + hip::Graph* g = reinterpret_cast(graph); + if (!hip::UserObject::isUserObjvalid(userObject) && !g->isUserObjGraphValid(userObject)) { HIP_RETURN(hipSuccess); } if (flags != hipGraphUserObjectMove) { @@ -2352,9 +2469,9 @@ hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, un } } else { //! if flag is UserObjMove delete userobj from list - hipUserObject::removeUSerObj(object); + hip::UserObject::removeUSerObj(userObject); } - graph->addUserObjGraph(object); + g->addUserObjGraph(userObject); HIP_RETURN(status); } @@ -2363,13 +2480,16 @@ hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, u if (graph == nullptr || object == nullptr || count == 0 || count > INT_MAX) { HIP_RETURN(hipErrorInvalidValue); } - if (!graph->isUserObjGraphValid(object) || object->referenceCount() < count) { + hip::UserObject* userObject = reinterpret_cast(object); + hip::Graph* g = reinterpret_cast(graph); + if (!g->isUserObjGraphValid(userObject) || userObject->referenceCount() < count) { HIP_RETURN(hipSuccess); } //! Obj is being destroyed - unsigned int releaseCount = (object->referenceCount() < count) ? object->referenceCount() : count; - if (object->referenceCount() == releaseCount) { - graph->RemoveUserObjGraph(object); + unsigned int releaseCount = + (userObject->referenceCount() < count) ? userObject->referenceCount() : count; + if (userObject->referenceCount() == releaseCount) { + g->RemoveUserObjGraph(userObject); } hipError_t status = hipUserObjectRelease(object, count); HIP_RETURN(status); @@ -2380,8 +2500,8 @@ hipError_t hipGraphKernelNodeCopyAttributes(hipGraphNode_t hSrc, hipGraphNode_t if (hSrc == nullptr || hDst == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(reinterpret_cast(hDst)->CopyAttr( - reinterpret_cast(hSrc))); + HIP_RETURN(reinterpret_cast(hDst)->CopyAttr( + reinterpret_cast(hSrc))); } hipError_t ihipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned int flags) { @@ -2395,7 +2515,7 @@ hipError_t ihipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned i return hipErrorOperatingSystem; } fout << "digraph dot {" << std::endl; - graph->GenerateDOT(fout, (hipGraphDebugDotFlags)flags); + reinterpret_cast(graph)->GenerateDOT(fout, (hipGraphDebugDotFlags)flags); fout << "}" << std::endl; fout.close(); return hipSuccess; @@ -2409,16 +2529,18 @@ hipError_t hipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned in hipError_t hipGraphNodeSetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, unsigned int isEnabled) { HIP_INIT_API(hipGraphNodeSetEnabled, hGraphExec, hNode, isEnabled); - if (hGraphExec == nullptr || hNode == nullptr || !hipGraphExec::isGraphExecValid(hGraphExec) || - !hipGraphNode::isNodeValid(hNode)) { + hip::GraphExec* graphExec = reinterpret_cast(hGraphExec); + hip::GraphNode* node = reinterpret_cast(hNode); + if (hGraphExec == nullptr || hNode == nullptr || !hip::GraphExec::isGraphExecValid(graphExec) || + !hip::GraphNode::isNodeValid(node)) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode); + hip::GraphNode* clonedNode = graphExec->GetClonedNode(node); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - if (!(hNode->GetType() == hipGraphNodeTypeKernel || hNode->GetType() == hipGraphNodeTypeMemcpy || - hNode->GetType() == hipGraphNodeTypeMemset)) { + if (!(node->GetType() == hipGraphNodeTypeKernel || node->GetType() == hipGraphNodeTypeMemcpy || + node->GetType() == hipGraphNodeTypeMemset)) { HIP_RETURN(hipErrorInvalidValue); } clonedNode->SetEnabled(isEnabled); @@ -2428,16 +2550,19 @@ hipError_t hipGraphNodeSetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNod hipError_t hipGraphNodeGetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, unsigned int* isEnabled) { HIP_INIT_API(hipGraphNodeGetEnabled, hGraphExec, hNode, isEnabled); + hip::GraphExec* graphExec = reinterpret_cast(hGraphExec); + hip::GraphNode* node = reinterpret_cast(hNode); + if (hGraphExec == nullptr || hNode == nullptr || isEnabled == nullptr || - !hipGraphExec::isGraphExecValid(hGraphExec) || !hipGraphNode::isNodeValid(hNode)) { + !hip::GraphExec::isGraphExecValid(graphExec) || !hip::GraphNode::isNodeValid(node)) { HIP_RETURN(hipErrorInvalidValue); } - hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode); + hip::GraphNode* clonedNode = graphExec->GetClonedNode(node); if (clonedNode == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - if (!(hNode->GetType() == hipGraphNodeTypeKernel || hNode->GetType() == hipGraphNodeTypeMemcpy || - hNode->GetType() == hipGraphNodeTypeMemset)) { + if (!(node->GetType() == hipGraphNodeTypeKernel || node->GetType() == hipGraphNodeTypeMemcpy || + node->GetType() == hipGraphNodeTypeMemset)) { HIP_RETURN(hipErrorInvalidValue); } *isEnabled = clonedNode->GetEnabled(); diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index cd8bf3c657..4d39c566d9 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -48,18 +48,19 @@ const char* GetGraphNodeTypeString(uint32_t op) { return case_string; }; -int hipGraphNode::nextID = 0; -int ihipGraph::nextID = 0; -std::unordered_set hipGraphNode::nodeSet_; -amd::Monitor hipGraphNode::nodeSetLock_{"Guards global node set"}; -std::unordered_set ihipGraph::graphSet_; -amd::Monitor ihipGraph::graphSetLock_{"Guards global graph set"}; -std::unordered_set hipGraphExec::graphExecSet_; -amd::Monitor hipGraphExec::graphExecSetLock_{"Guards global exec graph set"}; -std::unordered_set hipUserObject::ObjectSet_; -amd::Monitor hipUserObject::UserObjectLock_{"Guards global user object"}; +namespace hip { +int GraphNode::nextID = 0; +int Graph::nextID = 0; +std::unordered_set GraphNode::nodeSet_; +amd::Monitor GraphNode::nodeSetLock_{"Guards global node set"}; +std::unordered_set Graph::graphSet_; +amd::Monitor Graph::graphSetLock_{"Guards global graph set"}; +std::unordered_set GraphExec::graphExecSet_; +amd::Monitor GraphExec::graphExecSetLock_{"Guards global exec graph set"}; +std::unordered_set UserObject::ObjectSet_; +amd::Monitor UserObject::UserObjectLock_{"Guards global user object"}; -hipError_t hipGraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size_t count, +hipError_t GraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size_t count, hipMemcpyKind kind) { hipError_t status = ihipMemcpy_validate(dst, src, count, kind); if (status != hipSuccess) { @@ -100,7 +101,7 @@ hipError_t hipGraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size return hipSuccess; } -hipError_t hipGraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParams) { +hipError_t GraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParams) { hipError_t status = ihipMemcpy3D_validate(pNodeParams); if (status != hipSuccess) { return status; @@ -258,7 +259,7 @@ hipError_t hipGraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParam return hipSuccess; } -bool ihipGraph::isGraphValid(ihipGraph* pGraph) { +bool Graph::isGraphValid(Graph* pGraph) { amd::ScopedLock lock(graphSetLock_); if (graphSet_.find(pGraph) == graphSet_.end()) { return false; @@ -266,20 +267,20 @@ bool ihipGraph::isGraphValid(ihipGraph* pGraph) { return true; } -void ihipGraph::AddNode(const Node& node) { +void Graph::AddNode(const Node& node) { vertices_.emplace_back(node); ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add %s(%p)\n", GetGraphNodeTypeString(node->GetType()), node); node->SetParentGraph(this); } -void ihipGraph::RemoveNode(const Node& node) { +void Graph::RemoveNode(const Node& node) { vertices_.erase(std::remove(vertices_.begin(), vertices_.end(), node), vertices_.end()); delete node; } // root nodes are all vertices with 0 in-degrees -std::vector ihipGraph::GetRootNodes() const { +std::vector Graph::GetRootNodes() const { std::vector roots; for (auto entry : vertices_) { if (entry->GetInDegree() == 0) { @@ -293,7 +294,7 @@ std::vector ihipGraph::GetRootNodes() const { } // leaf nodes are all vertices with 0 out-degrees -std::vector ihipGraph::GetLeafNodes() const { +std::vector Graph::GetLeafNodes() const { std::vector leafNodes; for (auto entry : vertices_) { if (entry->GetOutDegree() == 0) { @@ -303,7 +304,7 @@ std::vector ihipGraph::GetLeafNodes() const { return leafNodes; } -size_t ihipGraph::GetLeafNodeCount() const { +size_t Graph::GetLeafNodeCount() const { int numLeafNodes = 0; for (auto entry : vertices_) { if (entry->GetOutDegree() == 0) { @@ -313,7 +314,7 @@ size_t ihipGraph::GetLeafNodeCount() const { return numLeafNodes; } -std::vector> ihipGraph::GetEdges() const { +std::vector> Graph::GetEdges() const { std::vector> edges; for (const auto& i : vertices_) { for (const auto& j : i->GetEdges()) { @@ -323,7 +324,7 @@ std::vector> ihipGraph::GetEdges() const { return edges; } -void ihipGraph::GetRunListUtil(Node v, std::unordered_map& visited, +void Graph::GetRunListUtil(Node v, std::unordered_map& visited, std::vector& singleList, std::vector>& parallelLists, std::unordered_map>& dependencies) { @@ -369,7 +370,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, +void Graph::GetRunList(std::vector>& parallelLists, std::unordered_map>& dependencies) { std::vector singleList; @@ -392,7 +393,7 @@ void ihipGraph::GetRunList(std::vector>& parallelLists, } } } -bool ihipGraph::TopologicalOrder(std::vector& TopoOrder) { +bool Graph::TopologicalOrder(std::vector& TopoOrder) { std::queue q; std::unordered_map inDegree; for (auto entry : vertices_) { @@ -418,10 +419,10 @@ bool ihipGraph::TopologicalOrder(std::vector& TopoOrder) { } return false; } -ihipGraph* ihipGraph::clone(std::unordered_map& clonedNodes) const { - ihipGraph* newGraph = new ihipGraph(device_, this); +Graph* Graph::clone(std::unordered_map& clonedNodes) const { + Graph* newGraph = new Graph(device_, this); for (auto entry : vertices_) { - hipGraphNode* node = entry->clone(); + GraphNode* node = entry->clone(); node->SetParentGraph(newGraph); newGraph->vertices_.push_back(node); clonedNodes[entry] = node; @@ -448,12 +449,12 @@ ihipGraph* ihipGraph::clone(std::unordered_map& clonedNodes) const { return newGraph; } -ihipGraph* ihipGraph::clone() const { +Graph* Graph::clone() const { std::unordered_map clonedNodes; return clone(clonedNodes); } -bool hipGraphExec::isGraphExecValid(hipGraphExec* pGraphExec) { +bool GraphExec::isGraphExecValid(GraphExec* pGraphExec) { amd::ScopedLock lock(graphExecSetLock_); if (graphExecSet_.find(pGraphExec) == graphExecSet_.end()) { return false; @@ -461,7 +462,7 @@ bool hipGraphExec::isGraphExecValid(hipGraphExec* pGraphExec) { return true; } -hipError_t hipGraphExec::CreateStreams(uint32_t num_streams) { +hipError_t GraphExec::CreateStreams(uint32_t num_streams) { parallel_streams_.reserve(num_streams); for (uint32_t i = 0; i < num_streams; ++i) { auto stream = new hip::Stream(hip::getCurrentDevice(), @@ -478,7 +479,7 @@ hipError_t hipGraphExec::CreateStreams(uint32_t num_streams) { return hipSuccess; } -hipError_t hipGraphExec::Init() { +hipError_t GraphExec::Init() { hipError_t status; size_t min_num_streams = 1; @@ -554,7 +555,7 @@ hipError_t FillCommands(std::vector>& parallelLists, } void UpdateStream(std::vector>& parallelLists, hip::Stream* stream, - hipGraphExec* ptr) { + GraphExec* ptr) { int i = 0; for (const auto& list : parallelLists) { // first parallel list will be launched on the same queue as parent @@ -572,7 +573,7 @@ void UpdateStream(std::vector>& parallelLists, hip::Stream* st } } -hipError_t hipGraphExec::Run(hipStream_t stream) { +hipError_t GraphExec::Run(hipStream_t stream) { hipError_t status; if (hip::getStream(stream) == nullptr) { @@ -590,7 +591,7 @@ hipError_t hipGraphExec::Run(hipStream_t stream) { if (repeatLaunch_ == true) { for (auto& node : topoOrder_) { if (node->GetType() == hipGraphNodeTypeMemAlloc && - static_cast(node)->IsActiveMem() == true) { + static_cast(node)->IsActiveMem() == true) { return hipErrorInvalidValue; } } @@ -633,3 +634,4 @@ hipError_t hipGraphExec::Run(hipStream_t stream) { ResetQueueIndex(); return status; } +} // namespace hip diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 709a1d8457..468c368c9d 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -34,28 +34,32 @@ #include "hip_platform.hpp" #include "hip_mempool_impl.hpp" #include "hip_vm.hpp" - -typedef hipGraphNode* Node; +namespace hip { +struct Graph; +struct GraphNode; +struct GraphExec; +struct UserObject; +typedef GraphNode* Node; hipError_t FillCommands(std::vector>& parallelLists, std::unordered_map>& nodeWaitLists, std::vector& topoOrder, std::vector& rootCommands, amd::Command*& endCommand, hip::Stream* stream); void UpdateStream(std::vector>& parallelLists, hip::Stream* stream, - hipGraphExec* ptr); + GraphExec* ptr); -struct hipUserObject : public amd::ReferenceCountedObject { +struct UserObject : public amd::ReferenceCountedObject { typedef void (*UserCallbackDestructor)(void* data); - static std::unordered_set ObjectSet_; + static std::unordered_set ObjectSet_; static amd::Monitor UserObjectLock_; public: - hipUserObject(UserCallbackDestructor callback, void* data, unsigned int flags) + UserObject(UserCallbackDestructor callback, void* data, unsigned int flags) : ReferenceCountedObject(), callback_(callback), data_(data), flags_(flags) { amd::ScopedLock lock(UserObjectLock_); ObjectSet_.insert(this); } - virtual ~hipUserObject() { + virtual ~UserObject() { amd::ScopedLock lock(UserObjectLock_); if (callback_ != nullptr) { callback_(data_); @@ -76,7 +80,7 @@ struct hipUserObject : public amd::ReferenceCountedObject { } } - static bool isUserObjvalid(hipUserObject* pUsertObj) { + static bool isUserObjvalid(UserObject* pUsertObj) { auto it = ObjectSet_.find(pUsertObj); if (it == ObjectSet_.end()) { return false; @@ -84,7 +88,7 @@ struct hipUserObject : public amd::ReferenceCountedObject { return true; } - static void removeUSerObj(hipUserObject* pUsertObj) { + static void removeUSerObj(UserObject* pUsertObj) { amd::ScopedLock lock(UserObjectLock_); auto it = ObjectSet_.find(pUsertObj); if (it != ObjectSet_.end()) { @@ -97,9 +101,9 @@ struct hipUserObject : public amd::ReferenceCountedObject { void* data_; unsigned int flags_; //! Disable default operator= - hipUserObject& operator=(const hipUserObject&) = delete; + UserObject& operator=(const UserObject&) = delete; //! Disable copy constructor - hipUserObject(const hipUserObject& obj) = delete; + UserObject(const UserObject& obj) = delete; }; struct hipGraphNodeDOTAttribute { @@ -154,7 +158,7 @@ struct hipGraphNodeDOTAttribute { } }; -struct hipGraphNode : public hipGraphNodeDOTAttribute { +struct GraphNode : public hipGraphNodeDOTAttribute { protected: hip::Stream* stream_ = nullptr; unsigned int id_; @@ -163,18 +167,18 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { std::vector edges_; std::vector dependencies_; bool visited_; - // count of in coming edges + // count of in coming edge s size_t inDegree_; // count of outgoing edges size_t outDegree_; static int nextID; - struct ihipGraph* parentGraph_; - static std::unordered_set nodeSet_; + struct Graph* parentGraph_; + static std::unordered_set nodeSet_; static amd::Monitor nodeSetLock_; unsigned int isEnabled_; public: - hipGraphNode(hipGraphNodeType type, std::string style = "", std::string shape = "", + GraphNode(hipGraphNodeType type, std::string style = "", std::string shape = "", std::string label = "") : type_(type), visited_(false), @@ -188,7 +192,7 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { nodeSet_.insert(this); } /// Copy Constructor - hipGraphNode(const hipGraphNode& node) : hipGraphNodeDOTAttribute(node) { + GraphNode(const GraphNode& node) : hipGraphNodeDOTAttribute(node) { type_ = node.type_; inDegree_ = node.inDegree_; outDegree_ = node.outDegree_; @@ -200,7 +204,7 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { isEnabled_ = node.isEnabled_; } - virtual ~hipGraphNode() { + virtual ~GraphNode() { for (auto node : edges_) { node->RemoveDependency(this); } @@ -212,7 +216,7 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { } // check node validity - static bool isNodeValid(hipGraphNode* pGraphNode) { + static bool isNodeValid(GraphNode* pGraphNode) { amd::ScopedLock lock(nodeSetLock_); if (pGraphNode == nullptr || nodeSet_.find(pGraphNode) == nodeSet_.end()) { return false; @@ -222,7 +226,7 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { hip::Stream* GetQueue() { return stream_; } - virtual void SetStream(hip::Stream* stream, hipGraphExec* ptr = nullptr) { + virtual void SetStream(hip::Stream* stream, GraphExec* ptr = nullptr) { stream_ = stream; } /// Create amd::command for the graph node @@ -238,7 +242,7 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { /// Returns graph node type hipGraphNodeType GetType() const { return type_; } /// Clone graph node - virtual hipGraphNode* clone() const = 0; + virtual GraphNode* clone() const = 0; /// Returns graph node indegree size_t GetInDegree() const { return inDegree_; } /// Updates indegree of the node @@ -331,10 +335,10 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { command->setBufferingState(true); } } - ihipGraph* GetParentGraph() { return parentGraph_; } - virtual ihipGraph* GetChildGraph() { return nullptr; } - void SetParentGraph(ihipGraph* graph) { parentGraph_ = graph; } - virtual hipError_t SetParams(hipGraphNode* node) { return hipSuccess; } + Graph* GetParentGraph() { return parentGraph_; } + virtual Graph* GetChildGraph() { return nullptr; } + void SetParentGraph(Graph* graph) { parentGraph_ = graph; } + virtual hipError_t SetParams(GraphNode* node) { return hipSuccess; } virtual void GenerateDOT(std::ostream& fout, hipGraphDebugDotFlags flag) {} virtual void GenerateDOTNode(size_t graphId, std::ostream& fout, hipGraphDebugDotFlags flag) { fout << "\n"; @@ -358,21 +362,21 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { void SetEnabled(unsigned int isEnabled) { isEnabled_ = isEnabled; } }; -struct ihipGraph { +struct Graph { std::vector vertices_; - const ihipGraph* pOriginalGraph_ = nullptr; - static std::unordered_set graphSet_; + const Graph* pOriginalGraph_ = nullptr; + static std::unordered_set graphSet_; static amd::Monitor graphSetLock_; - std::unordered_set graphUserObj_; + std::unordered_set graphUserObj_; unsigned int id_; static int nextID; hip::Device* device_; //!< HIP device object hip::MemoryPool* mem_pool_; //!< Memory pool, associated with this graph - std::unordered_set capturedNodes_; + std::unordered_set capturedNodes_; bool graphInstantiated_; public: - ihipGraph(hip::Device* device, const ihipGraph* original = nullptr) + Graph(hip::Device* device, const Graph* original = nullptr) : pOriginalGraph_(original) , id_(nextID++) , device_(device) { @@ -383,7 +387,7 @@ struct ihipGraph { graphInstantiated_ = false; } - ~ihipGraph() { + ~Graph() { for (auto node : vertices_) { delete node; } @@ -398,9 +402,9 @@ struct ihipGraph { } - void AddManualNodeDuringCapture(hipGraphNode* node) { capturedNodes_.insert(node); } + void AddManualNodeDuringCapture(GraphNode* node) { capturedNodes_.insert(node); } - std::unordered_set GetManualNodesDuringCapture() { return capturedNodes_; } + std::unordered_set GetManualNodesDuringCapture() { return capturedNodes_; } void RemoveManualNodesDuringCapture() { capturedNodes_.erase(capturedNodes_.begin(), capturedNodes_.end()); @@ -410,7 +414,7 @@ struct ihipGraph { int GetID() const { return id_; } // check graphs validity - static bool isGraphValid(ihipGraph* pGraph); + static bool isGraphValid(Graph* pGraph); /// add node to the graph void AddNode(const Node& node); @@ -428,21 +432,21 @@ struct ihipGraph { /// returns all the edges in the graph std::vector> GetEdges() const; // returns the original graph ptr if cloned - const ihipGraph* getOriginalGraph() const { return pOriginalGraph_; } + const Graph* getOriginalGraph() const { return pOriginalGraph_; } // Add user obj resource to graph - void addUserObjGraph(hipUserObject* pUserObj) { + void addUserObjGraph(UserObject* pUserObj) { amd::ScopedLock lock(graphSetLock_); graphUserObj_.insert(pUserObj); } // Check user obj resource from graph is valid - bool isUserObjGraphValid(hipUserObject* pUserObj) { + bool isUserObjGraphValid(UserObject* pUserObj) { if (graphUserObj_.find(pUserObj) == graphUserObj_.end()) { return false; } return true; } // Delete user obj resource from graph - void RemoveUserObjGraph(hipUserObject* pUserObj) { graphUserObj_.erase(pUserObj); } + void RemoveUserObjGraph(UserObject* pUserObj) { graphUserObj_.erase(pUserObj); } void GetRunListUtil(Node v, std::unordered_map& visited, std::vector& singleList, std::vector>& parallelLists, @@ -450,14 +454,14 @@ struct ihipGraph { void GetRunList(std::vector>& parallelLists, std::unordered_map>& dependencies); bool TopologicalOrder(std::vector& TopoOrder); - void GetUserObjs(std::unordered_set& graphExeUserObjs) { + void GetUserObjs(std::unordered_set& graphExeUserObjs) { for (auto userObj : graphUserObj_) { userObj->retain(); graphExeUserObjs.insert(userObj); } } - ihipGraph* clone(std::unordered_map& clonedNodes) const; - ihipGraph* clone() const; + Graph* clone(std::unordered_map& clonedNodes) const; + Graph* clone() const; void GenerateDOT(std::ostream& fout, hipGraphDebugDotFlags flag) { fout << "subgraph cluster_" << GetID() << " {" << std::endl; fout << "label=\"graph_" << GetID() <<"\"graph[style=\"dashed\"];\n"; @@ -543,7 +547,7 @@ struct ihipGraph { } }; -struct hipGraphExec { +struct GraphExec { std::vector> parallelLists_; // Topological order of the graph doesn't include nodes embedded as part of the child graph std::vector topoOrder_; @@ -552,16 +556,16 @@ struct hipGraphExec { uint currentQueueIndex_; std::unordered_map clonedNodes_; amd::Command* lastEnqueuedCommand_; - static std::unordered_set graphExecSet_; - std::unordered_set graphExeUserObj_; + static std::unordered_set graphExecSet_; + std::unordered_set graphExeUserObj_; static amd::Monitor graphExecSetLock_; uint64_t flags_ = 0; bool repeatLaunch_ = false; public: - hipGraphExec(std::vector& topoOrder, std::vector>& lists, + GraphExec(std::vector& topoOrder, std::vector>& lists, std::unordered_map>& nodeWaitLists, std::unordered_map& clonedNodes, - std::unordered_set& userObjs, + std::unordered_set& userObjs, uint64_t flags = 0) : parallelLists_(lists), topoOrder_(topoOrder), @@ -575,7 +579,7 @@ struct hipGraphExec { graphExecSet_.insert(this); } - ~hipGraphExec() { + ~GraphExec() { // new commands are launched for every launch they are destroyed as and when command is // terminated after it complete execution for (auto stream : parallel_streams_) { @@ -602,7 +606,7 @@ struct hipGraphExec { } // check executable graphs validity - static bool isGraphExecValid(hipGraphExec* pGraphExec); + static bool isGraphExecValid(GraphExec* pGraphExec); std::vector& GetNodes() { return topoOrder_; } @@ -613,30 +617,30 @@ struct hipGraphExec { hipError_t Run(hipStream_t stream); }; -struct hipChildGraphNode : public hipGraphNode { - struct ihipGraph* childGraph_; +struct ChildGraphNode : public GraphNode { + struct Graph* childGraph_; std::vector childGraphNodeOrder_; std::vector> parallelLists_; std::unordered_map> nodeWaitLists_; amd::Command* lastEnqueuedCommand_; public: - hipChildGraphNode(ihipGraph* g) : hipGraphNode(hipGraphNodeTypeGraph, "solid", "rectangle") { + ChildGraphNode(Graph* g) : GraphNode(hipGraphNodeTypeGraph, "solid", "rectangle") { childGraph_ = g->clone(); lastEnqueuedCommand_ = nullptr; } - ~hipChildGraphNode() { delete childGraph_; } + ~ChildGraphNode() { delete childGraph_; } - hipChildGraphNode(const hipChildGraphNode& rhs) : hipGraphNode(rhs) { + ChildGraphNode(const ChildGraphNode& rhs) : GraphNode(rhs) { childGraph_ = rhs.childGraph_->clone(); } - hipGraphNode* clone() const { - return new hipChildGraphNode(static_cast(*this)); + GraphNode* clone() const { + return new ChildGraphNode(static_cast(*this)); } - ihipGraph* GetChildGraph() { return childGraph_; } + Graph* GetChildGraph() { return childGraph_; } hipError_t GetNumParallelStreams(size_t &num) { if (false == TopologicalOrder(childGraphNodeOrder_)) { @@ -653,7 +657,7 @@ struct hipChildGraphNode : public hipGraphNode { return hipSuccess; } - void SetStream(hip::Stream* stream, hipGraphExec* ptr = nullptr) { + void SetStream(hip::Stream* stream, GraphExec* ptr = nullptr) { stream_ = stream; UpdateStream(parallelLists_, stream, ptr); } @@ -664,7 +668,7 @@ struct hipChildGraphNode : public hipGraphNode { // Create child graph node commands and set waitlists hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -709,7 +713,7 @@ struct hipChildGraphNode : public hipGraphNode { } } - hipError_t SetParams(const ihipGraph* childGraph) { + hipError_t SetParams(const Graph* childGraph) { const std::vector& newNodes = childGraph->GetNodes(); const std::vector& oldNodes = childGraph_->GetNodes(); for (std::vector::size_type i = 0; i != newNodes.size(); i++) { @@ -721,8 +725,8 @@ struct hipChildGraphNode : public hipGraphNode { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipChildGraphNode* childGraphNode = static_cast(node); + hipError_t SetParams(GraphNode* node) { + const ChildGraphNode* childGraphNode = static_cast(node); return SetParams(childGraphNode->childGraph_); } @@ -735,7 +739,7 @@ struct hipChildGraphNode : public hipGraphNode { } }; -class hipGraphKernelNode : public hipGraphNode { +class GraphKernelNode : public GraphNode { hipKernelNodeParams kernelParams_; unsigned int numParams_; hipKernelNodeAttrValue kernelAttr_; @@ -886,8 +890,8 @@ class hipGraphKernelNode : public hipGraphNode { return hipSuccess; } - hipGraphKernelNode(const hipKernelNodeParams* pNodeParams) - : hipGraphNode(hipGraphNodeTypeKernel, "bold", "octagon", "KERNEL") { + GraphKernelNode(const hipKernelNodeParams* pNodeParams) + : GraphNode(hipGraphNodeTypeKernel, "bold", "octagon", "KERNEL") { kernelParams_ = *pNodeParams; if (copyParams(pNodeParams) != hipSuccess) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to copy params"); @@ -896,7 +900,7 @@ class hipGraphKernelNode : public hipGraphNode { kernelAttrInUse_ = 0; } - ~hipGraphKernelNode() { freeParams(); } + ~GraphKernelNode() { freeParams(); } void freeParams() { // Deallocate memory allocated for kernargs passed via 'kernelParams' @@ -920,7 +924,7 @@ class hipGraphKernelNode : public hipGraphNode { } } - hipGraphKernelNode(const hipGraphKernelNode& rhs) : hipGraphNode(rhs) { + GraphKernelNode(const GraphKernelNode& rhs) : GraphNode(rhs) { kernelParams_ = rhs.kernelParams_; hipError_t status = copyParams(&rhs.kernelParams_); if (status != hipSuccess) { @@ -934,8 +938,8 @@ class hipGraphKernelNode : public hipGraphNode { } } - hipGraphNode* clone() const { - return new hipGraphKernelNode(static_cast(*this)); + GraphNode* clone() const { + return new GraphKernelNode(static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { @@ -945,7 +949,7 @@ class hipGraphKernelNode : public hipGraphNode { if (hipSuccess != status) { return status; } - status = hipGraphNode::CreateCommand(stream); + status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1033,7 +1037,7 @@ class hipGraphKernelNode : public hipGraphNode { } return hipSuccess; } - hipError_t CopyAttr(const hipGraphKernelNode* srcNode) { + hipError_t CopyAttr(const GraphKernelNode* srcNode) { if (kernelAttrInUse_ == 0 && srcNode->kernelAttrInUse_ == 0) { return hipSuccess; } @@ -1059,8 +1063,8 @@ class hipGraphKernelNode : public hipGraphNode { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphKernelNode* kernelNode = static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphKernelNode* kernelNode = static_cast(node); return SetParams(&kernelNode->kernelParams_); } @@ -1090,29 +1094,29 @@ class hipGraphKernelNode : public hipGraphNode { } }; -class hipGraphMemcpyNode : public hipGraphNode { +class GraphMemcpyNode : public GraphNode { hipMemcpy3DParms copyParams_; public: - hipGraphMemcpyNode(const hipMemcpy3DParms* pCopyParams) - : hipGraphNode(hipGraphNodeTypeMemcpy, "solid", "trapezium", "MEMCPY") { + GraphMemcpyNode(const hipMemcpy3DParms* pCopyParams) + : GraphNode(hipGraphNodeTypeMemcpy, "solid", "trapezium", "MEMCPY") { copyParams_ = *pCopyParams; } - ~hipGraphMemcpyNode() {} + ~GraphMemcpyNode() {} - hipGraphMemcpyNode(const hipGraphMemcpyNode& rhs) : hipGraphNode(rhs) { + GraphMemcpyNode(const GraphMemcpyNode& rhs) : GraphNode(rhs) { copyParams_ = rhs.copyParams_; } - hipGraphNode* clone() const { - return new hipGraphMemcpyNode(static_cast(*this)); + GraphNode* clone() const { + return new GraphMemcpyNode(static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { if (IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.kind)) { return hipSuccess; } - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1130,7 +1134,7 @@ class hipGraphMemcpyNode : public hipGraphNode { copyParams_.extent.depth, *hip::getStream(stream)); return; } - hipGraphNode::EnqueueCommands(stream); + GraphNode::EnqueueCommands(stream); } void GetParams(hipMemcpy3DParms* params) { @@ -1144,8 +1148,8 @@ class hipGraphMemcpyNode : public hipGraphNode { std::memcpy(©Params_, params, sizeof(hipMemcpy3DParms)); return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphMemcpyNode* memcpyNode = static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphMemcpyNode* memcpyNode = static_cast(node); return SetParams(&memcpyNode->copyParams_); } // ToDo: use this when commands are cloned and command params are to be updated @@ -1235,7 +1239,7 @@ class hipGraphMemcpyNode : public hipGraphNode { } }; -class hipGraphMemcpyNode1D : public hipGraphNode { +class GraphMemcpyNode1D : public GraphNode { protected: void* dst_; const void* src_; @@ -1243,25 +1247,25 @@ class hipGraphMemcpyNode1D : public hipGraphNode { hipMemcpyKind kind_; public: - hipGraphMemcpyNode1D(void* dst, const void* src, size_t count, hipMemcpyKind kind, + GraphMemcpyNode1D(void* dst, const void* src, size_t count, hipMemcpyKind kind, hipGraphNodeType type = hipGraphNodeTypeMemcpy) - : hipGraphNode(type, "solid", "trapezium", "MEMCPY"), + : GraphNode(type, "solid", "trapezium", "MEMCPY"), dst_(dst), src_(src), count_(count), kind_(kind) {} - ~hipGraphMemcpyNode1D() {} + ~GraphMemcpyNode1D() {} - hipGraphNode* clone() const { - return new hipGraphMemcpyNode1D(static_cast(*this)); + GraphNode* clone() const { + return new GraphMemcpyNode1D(static_cast(*this)); } virtual hipError_t CreateCommand(hip::Stream* stream) { if (IsHtoHMemcpy(dst_, src_, kind_)) { return hipSuccess; } - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1277,7 +1281,7 @@ class hipGraphMemcpyNode1D : public hipGraphNode { if (!isH2H) { if (commands_.empty()) return; // commands_ should have just 1 item - assert(commands_.size() == 1 && "Invalid command size in hipGraphMemcpyNode1D"); + assert(commands_.size() == 1 && "Invalid command size in GraphMemcpyNode1D"); } if (isEnabled_) { //HtoH @@ -1342,8 +1346,8 @@ class hipGraphMemcpyNode1D : public hipGraphNode { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphMemcpyNode1D* memcpy1DNode = static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphMemcpyNode1D* memcpy1DNode = static_cast(node); return SetParams(memcpy1DNode->dst_, memcpy1DNode->src_, memcpy1DNode->count_, memcpy1DNode->kind_); } @@ -1401,26 +1405,26 @@ class hipGraphMemcpyNode1D : public hipGraphNode { } }; -class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D { +class GraphMemcpyNodeFromSymbol : public GraphMemcpyNode1D { const void* symbol_; size_t offset_; public: - hipGraphMemcpyNodeFromSymbol(void* dst, const void* symbol, size_t count, size_t offset, + GraphMemcpyNodeFromSymbol(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind) - : hipGraphMemcpyNode1D(dst, nullptr, count, kind, hipGraphNodeTypeMemcpy), + : GraphMemcpyNode1D(dst, nullptr, count, kind, hipGraphNodeTypeMemcpy), symbol_(symbol), offset_(offset) {} - ~hipGraphMemcpyNodeFromSymbol() {} + ~GraphMemcpyNodeFromSymbol() {} - hipGraphNode* clone() const { - return new hipGraphMemcpyNodeFromSymbol( - static_cast(*this)); + GraphNode* clone() const { + return new GraphMemcpyNodeFromSymbol( + static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1486,32 +1490,32 @@ class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphMemcpyNodeFromSymbol* memcpyNode = - static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphMemcpyNodeFromSymbol* memcpyNode = + static_cast(node); return SetParams(memcpyNode->dst_, memcpyNode->symbol_, memcpyNode->count_, memcpyNode->offset_, memcpyNode->kind_); } }; -class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D { +class GraphMemcpyNodeToSymbol : public GraphMemcpyNode1D { const void* symbol_; size_t offset_; public: - hipGraphMemcpyNodeToSymbol(const void* symbol, const void* src, size_t count, size_t offset, + GraphMemcpyNodeToSymbol(const void* symbol, const void* src, size_t count, size_t offset, hipMemcpyKind kind) - : hipGraphMemcpyNode1D(nullptr, src, count, kind, hipGraphNodeTypeMemcpy), + : GraphMemcpyNode1D(nullptr, src, count, kind, hipGraphNodeTypeMemcpy), symbol_(symbol), offset_(offset) {} - ~hipGraphMemcpyNodeToSymbol() {} + ~GraphMemcpyNodeToSymbol() {} - hipGraphNode* clone() const { - return new hipGraphMemcpyNodeToSymbol(static_cast(*this)); + GraphNode* clone() const { + return new GraphMemcpyNodeToSymbol(static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1575,20 +1579,20 @@ class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphMemcpyNodeToSymbol* memcpyNode = - static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphMemcpyNodeToSymbol* memcpyNode = + static_cast(node); return SetParams(memcpyNode->src_, memcpyNode->symbol_, memcpyNode->count_, memcpyNode->offset_, memcpyNode->kind_); } }; -class hipGraphMemsetNode : public hipGraphNode { +class GraphMemsetNode : public GraphNode { hipMemsetParams memsetParams_; public: - hipGraphMemsetNode(const hipMemsetParams* pMemsetParams) - : hipGraphNode(hipGraphNodeTypeMemset, "solid", "invtrapezium", "MEMSET") { + GraphMemsetNode(const hipMemsetParams* pMemsetParams) + : GraphNode(hipGraphNodeTypeMemset, "solid", "invtrapezium", "MEMSET") { memsetParams_ = *pMemsetParams; size_t sizeBytes = 0; if (memsetParams_.height == 1) { @@ -1598,14 +1602,14 @@ class hipGraphMemsetNode : public hipGraphNode { } } - ~hipGraphMemsetNode() { } + ~GraphMemsetNode() { } // Copy constructor - hipGraphMemsetNode(const hipGraphMemsetNode& memsetNode) : hipGraphNode(memsetNode) { + GraphMemsetNode(const GraphMemsetNode& memsetNode) : GraphNode(memsetNode) { memsetParams_ = memsetNode.memsetParams_; } - hipGraphNode* clone() const { - return new hipGraphMemsetNode(static_cast(*this)); + GraphNode* clone() const { + return new GraphMemsetNode(static_cast(*this)); } std::string GetLabel(hipGraphDebugDotFlags flag) { @@ -1641,7 +1645,7 @@ class hipGraphMemsetNode : public hipGraphNode { } hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1726,27 +1730,27 @@ class hipGraphMemsetNode : public hipGraphNode { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphMemsetNode* memsetNode = static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphMemsetNode* memsetNode = static_cast(node); return SetParams(&memsetNode->memsetParams_); } }; -class hipGraphEventRecordNode : public hipGraphNode { +class GraphEventRecordNode : public GraphNode { hipEvent_t event_; public: - hipGraphEventRecordNode(hipEvent_t event) - : hipGraphNode(hipGraphNodeTypeEventRecord, "solid", "rectangle", "EVENT_RECORD"), + GraphEventRecordNode(hipEvent_t event) + : GraphNode(hipGraphNodeTypeEventRecord, "solid", "rectangle", "EVENT_RECORD"), event_(event) {} - ~hipGraphEventRecordNode() {} + ~GraphEventRecordNode() {} - hipGraphNode* clone() const { - return new hipGraphEventRecordNode(static_cast(*this)); + GraphNode* clone() const { + return new GraphEventRecordNode(static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1778,28 +1782,28 @@ class hipGraphEventRecordNode : public hipGraphNode { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphEventRecordNode* eventRecordNode = - static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphEventRecordNode* eventRecordNode = + static_cast(node); return SetParams(eventRecordNode->event_); } }; -class hipGraphEventWaitNode : public hipGraphNode { +class GraphEventWaitNode : public GraphNode { hipEvent_t event_; public: - hipGraphEventWaitNode(hipEvent_t event) - : hipGraphNode(hipGraphNodeTypeWaitEvent, "solid", "rectangle", "EVENT_WAIT"), + GraphEventWaitNode(hipEvent_t event) + : GraphNode(hipGraphNodeTypeWaitEvent, "solid", "rectangle", "EVENT_WAIT"), event_(event) {} - ~hipGraphEventWaitNode() {} + ~GraphEventWaitNode() {} - hipGraphNode* clone() const { - return new hipGraphEventWaitNode(static_cast(*this)); + GraphNode* clone() const { + return new GraphEventWaitNode(static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1831,32 +1835,32 @@ class hipGraphEventWaitNode : public hipGraphNode { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphEventWaitNode* eventWaitNode = static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphEventWaitNode* eventWaitNode = static_cast(node); return SetParams(eventWaitNode->event_); } }; -class hipGraphHostNode : public hipGraphNode { +class GraphHostNode : public GraphNode { hipHostNodeParams NodeParams_; public: - hipGraphHostNode(const hipHostNodeParams* NodeParams) - : hipGraphNode(hipGraphNodeTypeHost, "solid", "rectangle", "HOST") { + GraphHostNode(const hipHostNodeParams* NodeParams) + : GraphNode(hipGraphNodeTypeHost, "solid", "rectangle", "HOST") { NodeParams_ = *NodeParams; } - ~hipGraphHostNode() { } + ~GraphHostNode() { } - hipGraphHostNode(const hipGraphHostNode& hostNode) : hipGraphNode(hostNode) { + GraphHostNode(const GraphHostNode& hostNode) : GraphNode(hostNode) { NodeParams_ = hostNode.NodeParams_; } - hipGraphNode* clone() const { - return new hipGraphHostNode(static_cast(*this)); + GraphNode* clone() const { + return new GraphHostNode(static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1874,7 +1878,7 @@ class hipGraphHostNode : public hipGraphNode { void EnqueueCommands(hipStream_t stream) { if (!commands_.empty()) { - if (!commands_[0]->setCallback(CL_COMPLETE, hipGraphHostNode::Callback, &NodeParams_)) { + if (!commands_[0]->setCallback(CL_COMPLETE, GraphHostNode::Callback, &NodeParams_)) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed during setCallback"); } commands_[0]->enqueue(); @@ -1900,23 +1904,23 @@ class hipGraphHostNode : public hipGraphNode { return hipSuccess; } - hipError_t SetParams(hipGraphNode* node) { - const hipGraphHostNode* hostNode = static_cast(node); + hipError_t SetParams(GraphNode* node) { + const GraphHostNode* hostNode = static_cast(node); return SetParams(&hostNode->NodeParams_); } }; -class hipGraphEmptyNode : public hipGraphNode { +class GraphEmptyNode : public GraphNode { public: - hipGraphEmptyNode() : hipGraphNode(hipGraphNodeTypeEmpty, "solid", "rectangle", "EMPTY") {} - ~hipGraphEmptyNode() {} + GraphEmptyNode() : GraphNode(hipGraphNodeTypeEmpty, "solid", "rectangle", "EMPTY") {} + ~GraphEmptyNode() {} - hipGraphNode* clone() const { - return new hipGraphEmptyNode(static_cast(*this)); + GraphNode* clone() const { + return new GraphEmptyNode(static_cast(*this)); } hipError_t CreateCommand(hip::Stream* stream) { - hipError_t status = hipGraphNode::CreateCommand(stream); + hipError_t status = GraphNode::CreateCommand(stream); if (status != hipSuccess) { return status; } @@ -1929,7 +1933,7 @@ class hipGraphEmptyNode : public hipGraphNode { }; // ================================================================================================ -class hipGraphMemAllocNode : public hipGraphNode { +class GraphMemAllocNode : public GraphNode { hipMemAllocNodeParams node_params_; // Node parameters for memory allocation amd::Memory* va_ = nullptr; // Memory object, which holds a virtual address @@ -1938,7 +1942,7 @@ class hipGraphMemAllocNode : public hipGraphNode { class VirtualMemAllocNode : public amd::VirtualMapCommand { public: VirtualMemAllocNode(amd::HostQueue& queue, const amd::Event::EventWaitList& eventWaitList, - amd::Memory* va, size_t size, amd::Memory* memory, ihipGraph* graph) + amd::Memory* va, size_t size, amd::Memory* memory, Graph* graph) : VirtualMapCommand(queue, eventWaitList, va->getSvmPtr(), size, memory), va_(va), graph_(graph) {} @@ -1976,17 +1980,17 @@ class hipGraphMemAllocNode : public hipGraphNode { private: amd::Memory* va_; // Memory object with the new virtual address for mapping - ihipGraph* graph_; // Graph which allocates/maps memory + Graph* graph_; // Graph which allocates/maps memory }; public: - hipGraphMemAllocNode(const hipMemAllocNodeParams* node_params) - : hipGraphNode(hipGraphNodeTypeMemAlloc, "solid", "rectangle", "MEM_ALLOC") { + GraphMemAllocNode(const hipMemAllocNodeParams* node_params) + : GraphNode(hipGraphNodeTypeMemAlloc, "solid", "rectangle", "MEM_ALLOC") { node_params_ = *node_params; } - hipGraphMemAllocNode(const hipGraphMemAllocNode& rhs) - : hipGraphNode(rhs) { + GraphMemAllocNode(const GraphMemAllocNode& rhs) + : GraphNode(rhs) { node_params_ = rhs.node_params_; if (HIP_MEM_POOL_USE_VM) { assert(rhs.va_ != nullptr && "Graph MemAlloc runtime can't clone an invalid node!"); @@ -1995,18 +1999,18 @@ class hipGraphMemAllocNode : public hipGraphNode { } } - virtual ~hipGraphMemAllocNode() final { + virtual ~GraphMemAllocNode() final { if (va_ != nullptr) { va_->release(); } } - virtual hipGraphNode* clone() const final { - return new hipGraphMemAllocNode(static_cast(*this)); + virtual GraphNode* clone() const final { + return new GraphMemAllocNode(static_cast(*this)); } virtual hipError_t CreateCommand(hip::Stream* stream) final { - auto error = hipGraphNode::CreateCommand(stream); + auto error = GraphNode::CreateCommand(stream); if (!HIP_MEM_POOL_USE_VM) { auto ptr = Execute(stream_); } else { @@ -2077,14 +2081,14 @@ class hipGraphMemAllocNode : public hipGraphNode { }; // ================================================================================================ -class hipGraphMemFreeNode : public hipGraphNode { +class GraphMemFreeNode : public GraphNode { void* device_ptr_; // Device pointer of the freed memory // Derive the new class for VirtualMap command, since runtime has to free // real allocation after unmap is complete class VirtualMemFreeNode : public amd::VirtualMapCommand { public: - VirtualMemFreeNode(ihipGraph* graph, int device_id, amd::HostQueue& queue, + VirtualMemFreeNode(Graph* graph, int device_id, amd::HostQueue& queue, const amd::Event::EventWaitList& eventWaitList, void* ptr, size_t size, amd::Memory* memory) : VirtualMapCommand(queue, eventWaitList, ptr, size, memory) , graph_(graph), device_id_(device_id) {} @@ -2112,24 +2116,24 @@ class hipGraphMemFreeNode : public hipGraphNode { } private: - ihipGraph* graph_; // Graph, which has the execution of this command + Graph* graph_; // Graph, which has the execution of this command int device_id_; // Device ID where this command is executed }; public: - hipGraphMemFreeNode(void* dptr) - : hipGraphNode(hipGraphNodeTypeMemFree, "solid", "rectangle", "MEM_FREE") + GraphMemFreeNode(void* dptr) + : GraphNode(hipGraphNodeTypeMemFree, "solid", "rectangle", "MEM_FREE") , device_ptr_(dptr) {} - hipGraphMemFreeNode(const hipGraphMemFreeNode& rhs) : hipGraphNode(rhs) { + GraphMemFreeNode(const GraphMemFreeNode& rhs) : GraphNode(rhs) { device_ptr_ = rhs.device_ptr_; } - virtual hipGraphNode* clone() const final { - return new hipGraphMemFreeNode(static_cast(*this)); + virtual GraphNode* clone() const final { + return new GraphMemFreeNode(static_cast(*this)); } virtual hipError_t CreateCommand(hip::Stream* stream) final { - auto error = hipGraphNode::CreateCommand(stream); + auto error = GraphNode::CreateCommand(stream); if (!HIP_MEM_POOL_USE_VM) { Execute(stream_); } else { @@ -2159,3 +2163,4 @@ class hipGraphMemFreeNode : public hipGraphNode { *params = device_ptr_; } }; +} // namespace hip diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index 78ee55bc15..1cecf0547f 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -52,6 +52,13 @@ #define IHIP_IPC_MEM_HANDLE_SIZE 32 #define IHIP_IPC_MEM_RESERVED_SIZE LP64_SWITCH(20,12) +namespace hip { +struct Graph; +struct GraphNode; +struct GraphExec; +struct UserObject; +} + typedef struct ihipIpcMemHandle_st { char ipc_handle[IHIP_IPC_MEM_HANDLE_SIZE]; ///< ipc memory handle on ROCr size_t psize; @@ -238,7 +245,7 @@ namespace hip { /// Current capture status of the stream hipStreamCaptureStatus captureStatus_; /// Graph that is constructed with capture - hipGraph_t pCaptureGraph_; + hip::Graph* pCaptureGraph_; /// Based on mode stream capture places restrictions on API calls that can be made within or /// concurrently hipStreamCaptureMode captureMode_{hipStreamCaptureModeGlobal}; @@ -247,9 +254,9 @@ namespace hip { /// dependencies hipStream_t parentStream_ = nullptr; /// Last graph node captured in the stream - std::vector lastCapturedNodes_; + std::vector lastCapturedNodes_; /// dependencies removed via API hipStreamUpdateCaptureDependencies - std::vector removedDependencies_; + std::vector removedDependencies_; /// Derived streams/Paralell branches from the origin stream std::vector parallelCaptureStreams_; /// Capture events @@ -307,20 +314,20 @@ namespace hip { bool IsOriginStream() const { return originStream_; } void SetOriginStream() { originStream_ = true; } /// Returns captured graph - hipGraph_t GetCaptureGraph() const { return pCaptureGraph_; } + hip::Graph* GetCaptureGraph() const { return pCaptureGraph_; } /// Returns last captured graph node - const std::vector& GetLastCapturedNodes() const { return lastCapturedNodes_; } + const std::vector& GetLastCapturedNodes() const { return lastCapturedNodes_; } /// Set last captured graph node - void SetLastCapturedNode(hipGraphNode_t graphNode) { + void SetLastCapturedNode(hip::GraphNode* graphNode) { lastCapturedNodes_.clear(); lastCapturedNodes_.push_back(graphNode); } /// returns updated dependencies removed - const std::vector& GetRemovedDependencies() { + const std::vector& GetRemovedDependencies() { return removedDependencies_; } /// Append captured node via the wait event cross stream - void AddCrossCapturedNode(std::vector graphNodes, bool replace = false) { + void AddCrossCapturedNode(std::vector graphNodes, bool replace = false) { // replace dependencies as per flag hipStreamSetCaptureDependencies if (replace == true) { for (auto node : lastCapturedNodes_) { @@ -333,7 +340,7 @@ namespace hip { } } /// Set graph that is being captured - void SetCaptureGraph(hipGraph_t pGraph) { + void SetCaptureGraph(hip::Graph* pGraph) { pCaptureGraph_ = pGraph; captureStatus_ = hipStreamCaptureStatusActive; }