SWDEV-240806 - Added few Graph API Implementation

hipGraph APIs clone, childGraph, RemoveDependencies, GetEdges,GetDependencies, GetDependentNodes, GetType and DestroyNode
hipStream APIs GetCaptureInfo, GetCaptureInfo_v2, UpdateCaptureDependencies.

Change-Id: Ib0f4cb8ea335698eecdd6d744ffab1e954153673


[ROCm/clr commit: 1c51d82dd5]

projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h

@@ -974,6 +974,10 @@ typedef enum cudaStreamCaptureStatus hipStreamCaptureStatus;
 #define hipStreamCaptureStatusActive cudaStreamCaptureStatusActive
 #define hipStreamCaptureStatusInvalidated cudaStreamCaptureStatusInvalidated
 
+typedef enum cudaStreamUpdateCaptureDependenciesFlags hipStreamUpdateCaptureDependenciesFlags;
+#define hipStreamAddCaptureDependencies cudaStreamAddCaptureDependencies
+#define hipStreamSetCaptureDependencies cudaStreamSetCaptureDependencies
+
 /**
  * Stream CallBack struct
  */
@@ -2373,48 +2377,137 @@ inline static hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGra
       cudaGraphAddEmptyNode(pGraphNode, graph, pDependencies, numDependencies));
 }
 
-inline static hipError_t hipStreamWriteValue32(hipStream_t stream,
-                                               void* ptr, int32_t value, unsigned int flags) {
-   if (value < 0) {
-     printf("Warning! value is negative, CUDA accept positive values\n");
-   }
-   return hipCUResultTohipError(cuStreamWriteValue32(stream, reinterpret_cast<CUdeviceptr>(ptr),
-                                                             static_cast<cuuint32_t>(value), flags));
+inline static hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, int32_t value,
+                                               unsigned int flags) {
+    if (value < 0) {
+        printf("Warning! value is negative, CUDA accept positive values\n");
+    }
+    return hipCUResultTohipError(cuStreamWriteValue32(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                      static_cast<cuuint32_t>(value), flags));
 }
 
-inline static hipError_t hipStreamWriteValue64(hipStream_t stream,
-                                               void* ptr, int64_t value, unsigned int flags) {
-   if (value < 0) {
-     printf("Warning! value is negative, CUDA accept positive values\n");
-   }
-   return hipCUResultTohipError(cuStreamWriteValue64(stream, reinterpret_cast<CUdeviceptr>(ptr),
-                                                    static_cast<cuuint64_t>(value), flags));
+inline static hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, int64_t value,
+                                               unsigned int flags) {
+    if (value < 0) {
+        printf("Warning! value is negative, CUDA accept positive values\n");
+    }
+    return hipCUResultTohipError(cuStreamWriteValue64(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                      static_cast<cuuint64_t>(value), flags));
 }
 
-inline static hipError_t hipStreamWaitValue32(hipStream_t stream,
-                                              void* ptr, int32_t value, unsigned int flags,
+inline static hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, int32_t value,
+                                              unsigned int flags,
                                               uint32_t mask __dparm(0xFFFFFFFF)) {
-   if (value < 0) {
-     printf("Warning! value is negative, CUDA accept positive values\n");
-   }
-   if (mask != STREAM_OPS_WAIT_MASK_32) {
-     printf("Warning! mask will not have impact as CUDA ignores it.\n");
-   }
-   return hipCUResultTohipError(cuStreamWaitValue32(stream, reinterpret_cast<CUdeviceptr>(ptr),
-                                                    static_cast<cuuint32_t>(value), flags));
+    if (value < 0) {
+        printf("Warning! value is negative, CUDA accept positive values\n");
+    }
+    if (mask != STREAM_OPS_WAIT_MASK_32) {
+        printf("Warning! mask will not have impact as CUDA ignores it.\n");
+    }
+    return hipCUResultTohipError(cuStreamWaitValue32(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                     static_cast<cuuint32_t>(value), flags));
 }
 
-inline static hipError_t hipStreamWaitValue64(hipStream_t stream,
-                                              void* ptr, int64_t value, unsigned int flags,
+inline static hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, int64_t value,
+                                              unsigned int flags,
                                               uint64_t mask __dparm(0xFFFFFFFFFFFFFFFF)) {
-   if (value < 0) {
-     printf("Warning! value is negative, CUDA accept positive values\n");
-   }
-   if (mask != STREAM_OPS_WAIT_MASK_64) {
-     printf("Warning! mask will not have impact as CUDA ignores it.\n");
-   }
-   return hipCUResultTohipError(cuStreamWaitValue64(stream, reinterpret_cast<CUdeviceptr>(ptr),
-                                                    static_cast<cuuint64_t>(value), flags));
+    if (value < 0) {
+        printf("Warning! value is negative, CUDA accept positive values\n");
+    }
+    if (mask != STREAM_OPS_WAIT_MASK_64) {
+        printf("Warning! mask will not have impact as CUDA ignores it.\n");
+    }
+    return hipCUResultTohipError(cuStreamWaitValue64(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                     static_cast<cuuint64_t>(value), flags));
+}
+
+inline static hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from,
+                                                    const hipGraphNode_t* to,
+                                                    size_t numDependencies) {
+    return hipCUDAErrorTohipError(cudaGraphRemoveDependencies(graph, from, to, numDependencies));
+}
+
+
+inline static hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from,
+                                          hipGraphNode_t* to, size_t* numEdges) {
+    return hipCUDAErrorTohipError(cudaGraphGetEdges(graph, from, to, numEdges));
+}
+
+
+inline static hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node,
+                                                     hipGraphNode_t* pDependencies,
+                                                     size_t* pNumDependencies) {
+    return hipCUDAErrorTohipError(
+        cudaGraphNodeGetDependencies(node, pDependencies, pNumDependencies));
+}
+
+inline static hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node,
+                                                       hipGraphNode_t* pDependentNodes,
+                                                       size_t* pNumDependentNodes) {
+    return hipCUDAErrorTohipError(
+        cudaGraphNodeGetDependentNodes(node, pDependentNodes, pNumDependentNodes));
+}
+
+inline static hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType) {
+    return hipCUDAErrorTohipError(cudaGraphNodeGetType(node, pType));
+}
+
+inline static hipError_t hipGraphDestroyNode(hipGraphNode_t node) {
+    return hipCUDAErrorTohipError(cudaGraphDestroyNode(node));
+}
+
+inline static hipError_t hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph) {
+    return hipCUDAErrorTohipError(cudaGraphClone(pGraphClone, originalGraph));
+}
+
+inline static hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t originalNode,
+                                                 hipGraph_t clonedGraph) {
+    return hipCUDAErrorTohipError(cudaGraphNodeFindInClone(pNode, originalNode, clonedGraph));
+}
+
+inline static hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                                   const hipGraphNode_t* pDependencies,
+                                                   size_t numDependencies, hipGraph_t childGraph) {
+    return hipCUDAErrorTohipError(
+        cudaGraphAddChildGraphNode(pGraphNode, graph, pDependencies, numDependencies, childGraph));
+}
+
+inline static hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph) {
+    return hipCUDAErrorTohipError(cudaGraphChildGraphNodeGetGraph(node, pGraph));
+}
+
+inline static hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec,
+                                                             hipGraphNode_t node,
+                                                             hipGraph_t childGraph) {
+    return hipCUDAErrorTohipError(
+        cudaGraphExecChildGraphNodeSetParams(hGraphExec, node, childGraph));
+}
+
+
+inline static hipError_t hipStreamGetCaptureInfo(hipStream_t stream,
+                                                 hipStreamCaptureStatus* pCaptureStatus,
+                                                 unsigned long long* pId) {
+    return hipCUDAErrorTohipError(cudaStreamGetCaptureInfo(stream, pCaptureStatus, pId));
+}
+
+inline static hipError_t hipStreamGetCaptureInfo_v2(
+    hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
+    unsigned long long* id_out __dparm(0), hipGraph_t* graph_out __dparm(0),
+    const hipGraphNode_t** dependencies_out __dparm(0), size_t* numDependencies_out __dparm(0)) {
+    return hipCUDAErrorTohipError(cudaStreamGetCaptureInfo_v2(
+        stream, captureStatus_out, id_out, graph_out, dependencies_out, numDependencies_out));
+}
+
+inline static hipError_t hipStreamIsCapturing(hipStream_t stream,
+                                              hipStreamCaptureStatus* pCaptureStatus) {
+    return hipCUDAErrorTohipError(cudaStreamIsCapturing(stream, pCaptureStatus));
+}
+
+inline static hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream,
+                                                            hipGraphNode_t* dependencies,
+                                                            size_t numDependencies,
+                                                            unsigned int flags __dparm(0)) {
+    return hipCUDAErrorTohipError(cudaStreamUpdateCaptureDependencies(stream, dependencies, flags));
 }
 
 #endif  //__CUDACC__

projects/clr/hipamd/src/hip_graph.cpp

@@ -27,6 +27,14 @@
 thread_local std::vector<hipStream_t> g_captureStreams;
 std::unordered_map<amd::Command*, hipGraphExec_t> hipGraphExec::activeGraphExec_;
 
+inline void ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph,
+                            const hipGraphNode_t* pDependencies, size_t numDependencies) {
+  graph->AddNode(graphNode);
+  for (size_t i = 0; i < numDependencies; i++) {
+    pDependencies[i]->AddEdge(graphNode);
+  }
+}
+
 hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams) {
   hipFunction_t func = nullptr;
   hipError_t status =
@@ -70,14 +78,7 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
     return hipErrorInvalidDeviceFunction;
   }
   *pGraphNode = new hipGraphKernelNode(pNodeParams, func);
-  if (numDependencies == 0) {
-    graph->AddNode(*pGraphNode);
-  }
-  for (size_t i = 0; i < numDependencies; i++) {
-    if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) {
-      return hipErrorInvalidValue;
-    }
-  }
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
   return hipSuccess;
 }
 
@@ -88,16 +89,12 @@ hipError_t ihipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
       (numDependencies > 0 && pDependencies == nullptr) || pCopyParams == nullptr) {
     return hipErrorInvalidValue;
   }
-  ihipMemcpy3D_validate(pCopyParams);
+  hipError_t status = ihipMemcpy3D_validate(pCopyParams);
+  if (status != hipSuccess) {
+    return status;
+  }
   *pGraphNode = new hipGraphMemcpyNode(pCopyParams);
-  if (numDependencies == 0) {
-    graph->AddNode(*pGraphNode);
-  }
-  for (size_t i = 0; i < numDependencies; i++) {
-    if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) {
-      return hipErrorInvalidValue;
-    }
-  }
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
   return hipSuccess;
 }
 
@@ -108,16 +105,12 @@ hipError_t ihipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph
       (numDependencies > 0 && pDependencies == nullptr)) {
     return hipErrorInvalidValue;
   }
-  ihipMemcpy_validate(dst, src, count, kind);
+  hipError_t status = ihipMemcpy_validate(dst, src, count, kind);
+  if (status != hipSuccess) {
+    return status;
+  }
   *pGraphNode = new hipGraphMemcpyNode1D(dst, src, count, kind);
-  if (numDependencies == 0) {
-    graph->AddNode(*pGraphNode);
-  }
-  for (size_t i = 0; i < numDependencies; i++) {
-    if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) {
-      return hipErrorInvalidValue;
-    }
-  }
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
   return hipSuccess;
 }
 
@@ -128,25 +121,22 @@ hipError_t ihipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
       (numDependencies > 0 && pDependencies == nullptr) || pMemsetParams == nullptr) {
     return hipErrorInvalidValue;
   }
+  hipError_t status;
   if (pMemsetParams->height == 1) {
-    ihipMemset_validate(pMemsetParams->dst, pMemsetParams->value, pMemsetParams->elementSize,
-                        pMemsetParams->width * pMemsetParams->elementSize);
+    status =
+        ihipMemset_validate(pMemsetParams->dst, pMemsetParams->value, pMemsetParams->elementSize,
+                            pMemsetParams->width * pMemsetParams->elementSize);
   } else {
     auto sizeBytes = pMemsetParams->width * pMemsetParams->height * 1;
-    ihipMemset3D_validate(
+    status = ihipMemset3D_validate(
         {pMemsetParams->dst, pMemsetParams->pitch, pMemsetParams->width, pMemsetParams->height},
         pMemsetParams->value, {pMemsetParams->width, pMemsetParams->height, 1}, sizeBytes);
   }
-
+  if (status != hipSuccess) {
+    return status;
+  }
   *pGraphNode = new hipGraphMemsetNode(pMemsetParams);
-  if (numDependencies == 0) {
-    graph->AddNode(*pGraphNode);
-  }
-  for (size_t i = 0; i < numDependencies; i++) {
-    if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) {
-      return hipErrorInvalidValue;
-    }
-  }
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
   return hipSuccess;
 }
 
@@ -194,22 +184,17 @@ hipError_t capturehipMemcpyAsync(hipStream_t& stream, void*& dst, const void*& s
   ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy1D on stream : %p",
           stream);
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
-  hipGraphNode_t pGraphNode;
   hipGraph_t graph = nullptr;
   std::vector<hipGraphNode_t> pDependencies = s->GetLastCapturedNodes();
   size_t numDependencies = s->GetLastCapturedNodes().size();
   graph = s->GetCaptureGraph();
-  ihipMemcpy_validate(dst, src, sizeBytes, kind);
-  pGraphNode = new hipGraphMemcpyNode1D(dst, src, sizeBytes, kind);
-  if (numDependencies == 0) {
-    graph->AddNode(pGraphNode);
+  hipError_t status = ihipMemcpy_validate(dst, src, sizeBytes, kind);
+  if (status != hipSuccess) {
+    return status;
   }
-  for (size_t i = 0; i < numDependencies; i++) {
-    if (graph->AddEdge(pDependencies[i], pGraphNode) != hipSuccess) {
-      return hipErrorInvalidValue;
-    }
-  }
-  s->SetLastCapturedNode(pGraphNode);
+  hipGraphNode_t node = new hipGraphMemcpyNode1D(dst, src, sizeBytes, kind);
+  ihipGraphAddNode(node, graph, pDependencies.data(), numDependencies);
+  s->SetLastCapturedNode(node);
   return hipSuccess;
 }
 
@@ -318,13 +303,12 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns
   HIP_RETURN(hipSuccess);
 }
 
-hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus** pCaptureStatus) {
+hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) {
   HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus);
   if (stream == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  hipStreamCaptureStatus captureStatus = reinterpret_cast<hip::Stream*>(stream)->GetCaptureStatus();
-  *pCaptureStatus = &captureStatus;
+  *pCaptureStatus = reinterpret_cast<hip::Stream*>(stream)->GetCaptureStatus();
   HIP_RETURN(hipSuccess);
 }
 
@@ -363,8 +347,22 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) {
     HIP_RETURN(hipErrorStreamCaptureInvalidated);
   }
   // check if all parallel streams have joined
+  // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do
+  // not result in hipErrorStreamCaptureUnjoined
   if (s->GetCaptureGraph()->GetLeafNodeCount() != 1) {
-    return hipErrorStreamCaptureUnjoined;
+    std::vector<hipGraphNode_t> leafNodes = s->GetCaptureGraph()->GetLeafNodes();
+    const std::vector<hipGraphNode_t>& removedDepNodes = s->GetRemovedDependencies();
+    bool foundInRemovedDep = false;
+    for (auto leafNode : leafNodes) {
+      for (auto node : removedDepNodes) {
+        if (node == leafNode) {
+          foundInRemovedDep = true;
+        }
+      }
+    }
+    if (foundInRemovedDep == false) {
+      return hipErrorStreamCaptureUnjoined;
+    }
   }
   *pGraph = s->GetCaptureGraph();
   // end capture on all streams/events part of graph capture
@@ -379,6 +377,9 @@ hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags) {
 
 hipError_t hipGraphDestroy(hipGraph_t graph) {
   HIP_INIT_API(hipGraphDestroy, graph);
+  if (graph == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
   delete graph;
   HIP_RETURN(hipSuccess);
 }
@@ -430,15 +431,21 @@ hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
     HIP_RETURN(hipErrorInvalidValue);
   }
   *pGraphNode = new hipGraphEmptyNode();
-  if (numDependencies == 0) {
-    graph->AddNode(*pGraphNode);
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                     const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                     hipGraph_t childGraph) {
+  HIP_INIT_API(hipGraphAddChildGraphNode, pGraphNode, pDependencies, numDependencies, childGraph);
+  if (pGraphNode == nullptr || graph == nullptr ||
+      (numDependencies > 0 && pDependencies == nullptr) || childGraph == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
   }
-  for (size_t i = 0; i < numDependencies; i++) {
-    if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) {
-      return hipErrorInvalidValue;
-    }
-  }
-  return hipSuccess;
+  *pGraphNode = new hipChildGraphNode(childGraph);
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
+  HIP_RETURN(hipSuccess);
 }
 
 hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
@@ -464,6 +471,9 @@ hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
 
 hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) {
   HIP_INIT_API(hipGraphExecDestroy, pGraphExec);
+  if (pGraphExec == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
   delete pGraphExec;
   HIP_RETURN(hipSuccess);
 }
@@ -474,6 +484,9 @@ hipError_t ihipGraphlaunch(hipGraphExec_t graphExec, hipStream_t stream) {
 
 hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) {
   HIP_INIT_API(hipGraphLaunch, graphExec, stream);
+  if (graphExec == nullptr || stream == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
   HIP_RETURN_DURATION(ihipGraphlaunch(graphExec, stream));
 }
 
@@ -482,22 +495,43 @@ hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* num
   if (graph == nullptr || numNodes == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  *numNodes = graph->GetNodeCount();
-  if (*numNodes > 0) {
-    nodes = graph->GetNodes().data();
+  const std::vector<hipGraphNode_t>& graphNodes = graph->GetNodes();
+  if (nodes == nullptr) {
+    *numNodes = graphNodes.size();
+    HIP_RETURN(hipSuccess);
+  } else if (*numNodes < graphNodes.size()) {
+    HIP_RETURN(hipErrorInvalidValue);
   }
-  HIP_RETURN(hipSuccess);
+  for (int i = 0; i < graphNodes.size(); i++) {
+    nodes[i] = graphNodes[i];
+  }
+  for (int i = graphNodes.size(); i < *numNodes; i++) {
+    nodes[i] = nullptr;
+  }
+  *numNodes = graphNodes.size();
 }
 
 hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes,
                                 size_t* pNumRootNodes) {
   HIP_INIT_API(hipGraphGetRootNodes, graph, pRootNodes, pNumRootNodes);
+
   if (graph == nullptr || pNumRootNodes == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  std::vector<Node> rootNodes = graph->GetRootNodes();
-  pRootNodes = rootNodes.data();
-  *pNumRootNodes = rootNodes.size();
+  const std::vector<hipGraphNode_t> nodes = graph->GetRootNodes();
+  if (pRootNodes == nullptr) {
+    *pNumRootNodes = nodes.size();
+    HIP_RETURN(hipSuccess);
+  } else if (*pNumRootNodes < nodes.size()) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  for (int i = 0; i < nodes.size(); i++) {
+    pRootNodes[i] = nodes[i];
+  }
+  for (int i = nodes.size(); i < *pNumRootNodes; i++) {
+    pRootNodes[i] = nullptr;
+  }
+  *pNumRootNodes = nodes.size();
   HIP_RETURN(hipSuccess);
 }
 
@@ -568,9 +602,7 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
     HIP_RETURN(hipErrorInvalidValue);
   }
   for (size_t i = 0; i < numDependencies; i++) {
-    if (graph->AddEdge(from[i], to[i]) != hipSuccess) {
-      HIP_RETURN(hipErrorInvalidValue);
-    }
+    from[i]->AddEdge(to[i]);
   }
   HIP_RETURN(hipSuccess);
 }
@@ -583,3 +615,220 @@ hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo
   }
   return reinterpret_cast<hipGraphKernelNode*>(node)->SetCommandParams(pNodeParams);
 }
+
+hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph) {
+  HIP_INIT_API(hipGraphChildGraphNodeGetGraph, node, pGraph);
+  if (node == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  *pGraph = reinterpret_cast<hipChildGraphNode*>(node)->GetChildGraph();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
+                                               hipGraph_t childGraph) {
+  HIP_INIT_API(hipGraphExecChildGraphNodeSetParams, hGraphExec, node, childGraph);
+  if (hGraphExec == nullptr || node == nullptr || childGraph == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
+                                   unsigned long long* pId) {
+  HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId);
+  if (pCaptureStatus == nullptr || pId == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  if (stream == nullptr) {
+    HIP_RETURN(hipErrorStreamCaptureImplicit);
+  }
+  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+  *pCaptureStatus = s->GetCaptureStatus();
+  *pId = s->GetCaptureID();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
+                                      unsigned long long* id_out, hipGraph_t* graph_out,
+                                      const hipGraphNode_t** dependencies_out,
+                                      size_t* numDependencies_out) {
+  HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out,
+               dependencies_out, numDependencies_out);
+  if (stream == nullptr) {
+    HIP_RETURN(hipErrorStreamCaptureImplicit);
+  }
+  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+  *captureStatus_out = s->GetCaptureStatus();
+  if (*captureStatus_out == hipStreamCaptureStatusActive) {
+    if (id_out != nullptr) {
+      *id_out = s->GetCaptureID();
+    }
+    if (graph_out != nullptr) {
+      *graph_out = s->GetCaptureGraph();
+    }
+    if (dependencies_out != nullptr && numDependencies_out != nullptr) {
+      *dependencies_out = s->GetLastCapturedNodes().data();
+      *numDependencies_out = s->GetLastCapturedNodes().size();
+    }
+  }
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies,
+                                              size_t numDependencies, unsigned int flags) {
+  HIP_INIT_API(hipStreamUpdateCaptureDependencies, stream, dependencies, numDependencies, flags);
+  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+  if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) {
+    return hipErrorIllegalState;
+  }
+  if (numDependencies > 0 && dependencies == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  std::vector<hipGraphNode_t> depNodes;
+  for (int i = 0; i < numDependencies; i++) {
+    depNodes.push_back(dependencies[i]);
+  }
+  if (flags == hipStreamAddCaptureDependencies) {
+    s->AddCrossCapturedNode(depNodes);
+  } else if (flags == hipStreamSetCaptureDependencies) {
+    bool replace = true;
+    s->AddCrossCapturedNode(depNodes, replace);
+  }
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from,
+                                      const hipGraphNode_t* to, size_t numDependencies) {
+  HIP_INIT_API(hipGraphRemoveDependencies, graph, from, to, numDependencies);
+  if (graph == nullptr || from == nullptr || to == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  for (size_t i = 0; i < numDependencies; i++) {
+    from[i]->RemoveEdge(to[i]);
+  }
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode_t* to,
+                            size_t* numEdges) {
+  HIP_INIT_API(hipGraphGetEdges, graph, from, to, numEdges);
+  if (graph == nullptr || numEdges == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  const std::vector<std::pair<Node, Node>> edges = graph->GetEdges();
+  // returns only the number of edges in numEdges when from and to are null
+  if (from == nullptr && to == nullptr) {
+    *numEdges = edges.size();
+    return hipSuccess;
+  } else if (*numEdges < edges.size()) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  for (int i = 0; i < edges.size(); i++) {
+    from[i] = edges[i].first;
+    to[i] = edges[i].second;
+  }
+  // If numEdges > actual number of edges, the remaining entries in from and to will be set to NULL
+  for (int i = edges.size(); i < *numEdges; i++) {
+    from[i] = nullptr;
+    to[i] = nullptr;
+  }
+  *numEdges = edges.size();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDependencies,
+                                       size_t* pNumDependencies) {
+  HIP_INIT_API(hipGraphNodeGetDependencies, node, pDependencies, pNumDependencies);
+  if (node == nullptr || pNumDependencies == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  const std::vector<hipGraphNode_t>& dependencies = node->GetDependencies();
+  if (pDependencies == NULL) {
+    *pNumDependencies = dependencies.size();
+    HIP_RETURN(hipSuccess);
+  } else if (*pNumDependencies < dependencies.size()) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  for (int i = 0; i < dependencies.size(); i++) {
+    pDependencies[i] = dependencies[i];
+  }
+  // pNumDependencies > actual number of dependencies, the remaining entries in pDependencies will
+  // be set to NULL
+  for (int i = dependencies.size(); i < *pNumDependencies; i++) {
+    pDependencies[i] = nullptr;
+  }
+  *pNumDependencies = dependencies.size();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pDependentNodes,
+                                         size_t* pNumDependentNodes) {
+  HIP_INIT_API(hipGraphNodeGetDependentNodes, node, pDependentNodes, pNumDependentNodes);
+  if (node == nullptr || pNumDependentNodes == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  const std::vector<hipGraphNode_t>& dependents = node->GetEdges();
+  if (pDependentNodes == NULL) {
+    *pNumDependentNodes = dependents.size();
+    HIP_RETURN(hipSuccess);
+  } else if (*pNumDependentNodes < dependents.size()) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  for (int i = 0; i < dependents.size(); i++) {
+    pDependentNodes[i] = dependents[i];
+  }
+  // pNumDependentNodes > actual number of dependents, the remaining entries in pDependentNodes will
+  // be set to NULL
+  for (int i = dependents.size(); i < *pNumDependentNodes; i++) {
+    pDependentNodes[i] = nullptr;
+  }
+  *pNumDependentNodes = dependents.size();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType) {
+  HIP_INIT_API(hipGraphNodeGetType, node, pType);
+  if (node == nullptr || pType == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  *pType = node->GetType();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphDestroyNode(hipGraphNode_t node) {
+  HIP_INIT_API(hipGraphDestroyNode, node);
+  if (node == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  node->GetParentGraph()->RemoveNode(node);
+  // Takescare of removing its dependencies and dependent nodes
+  delete node;
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph) {
+  HIP_INIT_API(hipGraphClone, pGraphClone, originalGraph);
+  if (originalGraph == nullptr || pGraphClone == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  *pGraphClone = originalGraph->clone();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t originalNode,
+                                   hipGraph_t clonedGraph) {
+  HIP_INIT_API(hipGraphNodeFindInClone, pNode, originalNode, clonedGraph);
+  if (pNode == nullptr || originalNode == nullptr || clonedGraph == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  for (auto node : clonedGraph->GetNodes()) {
+    if (node->GetID() == originalNode->GetID()) {
+      *pNode = node;
+      HIP_RETURN(hipSuccess);
+    }
+  }
+  HIP_RETURN(hipErrorInvalidValue);
+}

projects/clr/hipamd/src/hip_graph_internal.cpp

@@ -45,49 +45,24 @@ const char* GetGraphNodeTypeString(uint32_t op) {
   return case_string;
 };
 
-hipError_t ihipGraph::AddNode(const Node& node) {
+int hipGraphNode::nextID = 0;
+
+void ihipGraph::AddNode(const Node& node) {
   vertices_.emplace_back(node);
-  nodeOutDegree_[node] = 0;
-  nodeInDegree_[node] = 0;
-  node->SetLevel(0);
   ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add %s(%p)\n",
           GetGraphNodeTypeString(node->GetType()), node);
-  return hipSuccess;
+  node->SetParentGraph(this);
 }
 
-hipError_t ihipGraph::AddEdge(const Node& parentNode, const Node& childNode) {
-  // if vertice doesn't exist, add it to the graph
-  if (std::find(vertices_.begin(), vertices_.end(), parentNode) == vertices_.end()) {
-    AddNode(parentNode);
-  }
-  if (std::find(vertices_.begin(), vertices_.end(), childNode) == vertices_.end()) {
-    AddNode(childNode);
-  }
-  // Check if edge already exists
-  auto connectedEdges = edges_.find(parentNode);
-  if (connectedEdges != edges_.end()) {
-    if (std::find(connectedEdges->second.begin(), connectedEdges->second.end(), childNode) !=
-        connectedEdges->second.end()) {
-      return hipSuccess;
-    }
-    connectedEdges->second.emplace_back(childNode);
-  } else {
-    edges_[parentNode] = {childNode};
-  }
-  nodeOutDegree_[parentNode]++;
-  nodeInDegree_[childNode]++;
-  childNode->SetLevel(std::max(childNode->GetLevel(), parentNode->GetLevel() + 1));
-  ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add edge btwn %s(%p) - %s(%p)\n",
-          GetGraphNodeTypeString(parentNode->GetType()), parentNode,
-          GetGraphNodeTypeString(childNode->GetType()), childNode);
-  return hipSuccess;
+void ihipGraph::RemoveNode(const Node& node) {
+  vertices_.erase(std::remove(vertices_.begin(), vertices_.end(), node), vertices_.end());
 }
 
 // root nodes are all vertices with 0 in-degrees
 std::vector<Node> ihipGraph::GetRootNodes() const {
   std::vector<Node> roots;
   for (auto entry : vertices_) {
-    if (nodeInDegree_.at(entry) == 0) {
+    if (entry->GetInDegree() == 0) {
       roots.push_back(entry);
       ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] root node: %s(%p)\n",
               GetGraphNodeTypeString(entry->GetType()), entry);
@@ -101,7 +76,7 @@ std::vector<Node> ihipGraph::GetRootNodes() const {
 std::vector<Node> ihipGraph::GetLeafNodes() const {
   std::vector<Node> leafNodes;
   for (auto entry : vertices_) {
-    if (nodeOutDegree_.at(entry) == 0) {
+    if (entry->GetOutDegree() == 0) {
       leafNodes.push_back(entry);
     }
   }
@@ -111,7 +86,7 @@ std::vector<Node> ihipGraph::GetLeafNodes() const {
 size_t ihipGraph::GetLeafNodeCount() const {
   int numLeafNodes = 0;
   for (auto entry : vertices_) {
-    if (nodeOutDegree_.at(entry) == 0) {
+    if (entry->GetOutDegree() == 0) {
       numLeafNodes++;
     }
   }
@@ -120,23 +95,23 @@ size_t ihipGraph::GetLeafNodeCount() const {
 
 std::vector<std::pair<Node, Node>> ihipGraph::GetEdges() const {
   std::vector<std::pair<Node, Node>> edges;
-  for (const auto& i : edges_) {
-    for (const auto& j : i.second) {
-      edges.push_back(std::make_pair(i.first, j));
+  for (const auto& i : vertices_) {
+    for (const auto& j : i->GetEdges()) {
+      edges.push_back(std::make_pair(i, j));
     }
   }
   return edges;
 }
 
 void ihipGraph::GetRunListUtil(Node v, std::unordered_map<Node, bool>& visited,
-                              std::vector<Node>& singleList,
-                              std::vector<std::vector<Node>>& parallelLists,
-                              std::unordered_map<Node, std::vector<Node>>& dependencies) {
+                               std::vector<Node>& singleList,
+                               std::vector<std::vector<Node>>& parallelLists,
+                               std::unordered_map<Node, std::vector<Node>>& dependencies) {
   // Mark the current node as visited.
   visited[v] = true;
   singleList.push_back(v);
   // Recurse for all the vertices adjacent to this vertex
-  for (auto& adjNode : edges_[v]) {
+  for (auto& adjNode : v->GetEdges()) {
     if (!visited[adjNode]) {
       // For the parallel list nodes add parent as the dependency
       if (singleList.empty()) {
@@ -175,7 +150,7 @@ void ihipGraph::GetRunListUtil(Node v, std::unordered_map<Node, bool>& visited,
 // The function to do Topological Sort.
 // It uses recursive GetRunListUtil()
 void ihipGraph::GetRunList(std::vector<std::vector<Node>>& parallelLists,
-                          std::unordered_map<Node, std::vector<Node>>& dependencies) {
+                           std::unordered_map<Node, std::vector<Node>>& dependencies) {
   std::vector<Node> singleList;
 
   // Mark all the vertices as not visited
@@ -184,6 +159,8 @@ void ihipGraph::GetRunList(std::vector<std::vector<Node>>& parallelLists,
 
   // Call the recursive helper function for all vertices one by one
   for (auto node : vertices_) {
+    // If the node has embedded child graph
+    node->GetRunList(parallelLists, dependencies);
     if (visited[node] == false) {
       GetRunListUtil(node, visited, singleList, parallelLists, dependencies);
     }
@@ -196,7 +173,7 @@ void ihipGraph::GetRunList(std::vector<std::vector<Node>>& parallelLists,
   }
 }
 
-hipError_t ihipGraph::LevelOrder(std::vector<Node>& levelOrder) {
+void ihipGraph::LevelOrder(std::vector<Node>& levelOrder) {
   std::vector<Node> roots = GetRootNodes();
   std::unordered_map<Node, bool> visited;
   std::queue<Node> q;
@@ -209,7 +186,7 @@ hipError_t ihipGraph::LevelOrder(std::vector<Node>& levelOrder) {
     Node& node = q.front();
     q.pop();
     levelOrder.push_back(node);
-    for (const auto& i : edges_[node]) {
+    for (const auto& i : node->GetEdges()) {
       if (visited.find(i) == visited.end() && i->GetLevel() == (node->GetLevel() + 1)) {
         q.push(i);
         ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] %s(%p) level:%d \n",
@@ -218,7 +195,36 @@ hipError_t ihipGraph::LevelOrder(std::vector<Node>& levelOrder) {
       }
     }
   }
-  return hipSuccess;
+}
+
+ihipGraph* ihipGraph::clone() const {
+  ihipGraph* newGraph = new ihipGraph();
+  std::unordered_map<Node, Node> clonedNodes;
+  for (auto entry : vertices_) {
+    hipGraphNode* node = entry->clone();
+    newGraph->vertices_.push_back(node);
+    clonedNodes[entry] = node;
+  }
+  std::vector<Node> dependancies;
+  std::vector<Node> clonedEdges;
+  std::vector<Node> clonedDependencies;
+  for (auto node : vertices_) {
+    const std::vector<Node>& edges = node->GetEdges();
+    clonedEdges.clear();
+    for (auto edge : edges) {
+      clonedEdges.push_back(clonedNodes[edge]);
+    }
+    clonedNodes[node]->SetEdges(clonedEdges);
+  }
+  for (auto node : vertices_) {
+    const std::vector<Node>& dependencies = node->GetDependencies();
+    clonedDependencies.clear();
+    for (auto dep : dependencies) {
+      clonedDependencies.push_back(clonedNodes[dep]);
+    }
+    clonedNodes[node]->SetDependencies(clonedDependencies);
+  }
+  return newGraph;
 }
 
 hipError_t hipGraphExec::CreateQueues() {
@@ -253,25 +259,31 @@ hipError_t hipGraphExec::FillCommands() {
     }
     i++;
   }
+
+  i = 0;
+  // For nodes that has embedded child graph
+  for (const auto& list : parallelLists_) {
+    for (auto& node : list) {
+      node->UpdateEventWaitLists();
+    }
+    i++;
+  }
+
   // Add waitlists for all the commands
-  for (auto& node : levelOrder_) {
-    auto nodeWaitList = nodeWaitLists_.find(node);
-    if (nodeWaitList != nodeWaitLists_.end()) {
-      amd::Command::EventWaitList waitList;
-      for (auto depNode : nodeWaitList->second) {
-        for (auto command : depNode->GetCommands()) {
-          waitList.push_back(command);
-        }
-      }
-      for (auto command : nodeWaitList->first->GetCommands()) {
-        command->updateEventWaitList(waitList);
+  for (auto entry : nodeWaitLists_) {
+    amd::Command::EventWaitList waitList;
+    for (auto depNode : entry.second) {
+      for (auto command : depNode->GetCommands()) {
+        waitList.push_back(command);
       }
     }
+    for (auto command : entry.first->GetCommands()) {
+      command->updateEventWaitList(waitList);
+    }
   }
   return status;
 }
 
-
 hipError_t hipGraphExec::Init() {
   hipError_t status;
   status = CreateQueues();
@@ -297,9 +309,7 @@ void hipGraphExec::ResetGraph(cl_event event, cl_int command_exec_status, void*
       hipGraphExec::activeGraphExec_[reinterpret_cast<amd::Command*>(user_data)];
   if (graphExec != nullptr) {
     for (auto& node : graphExec->levelOrder_) {
-      for (auto& command : node->GetCommands()) {
-        command->resetStatus(CL_INT_MAX);
-      }
+      node->ResetStatus();
     }
     graphExec->rootCommand_->resetStatus(CL_INT_MAX);
     graphExec->bExecPending_.store(false);
@@ -308,7 +318,7 @@ void hipGraphExec::ResetGraph(cl_event event, cl_int command_exec_status, void*
   }
 }
 
-hipError_t hipGraphExec::UpdateGraphToWaitOnRoot() {
+void hipGraphExec::UpdateGraphToWaitOnRoot() {
   for (auto& singleList : parallelLists_) {
     amd::Command::EventWaitList waitList;
     waitList.push_back(rootCommand_);
@@ -319,7 +329,6 @@ hipError_t hipGraphExec::UpdateGraphToWaitOnRoot() {
       }
     }
   }
-  return hipSuccess;
 }
 
 hipError_t hipGraphExec::Run(hipStream_t stream) {
@@ -342,9 +351,7 @@ hipError_t hipGraphExec::Run(hipStream_t stream) {
   }
   rootCommand_->enqueue();
   for (auto& node : levelOrder_) {
-    for (auto& command : node->GetCommands()) {
-      command->enqueue();
-    }
+    node->EnqueueCommands();
   }
 
   amd::Command* command = new amd::Marker(*queue, false, graphLastCmdWaitList_);
@@ -362,7 +369,7 @@ hipError_t hipGraphExec::Run(hipStream_t stream) {
   command->enqueue();
 
   // Add the new barrier to stall the stream, until the callback is done
-  amd::Command::EventWaitList eventWaitList;;
+  amd::Command::EventWaitList eventWaitList;
   eventWaitList.push_back(command);
   amd::Command* block_command = new amd::Marker(*queue, !kMarkerDisableFlush, eventWaitList);
   if (block_command == nullptr) {

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -33,45 +33,161 @@
 
 typedef hipGraphNode* Node;
 hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams);
-
 struct hipGraphNode {
  protected:
   uint32_t level_;
+  unsigned int id_;
   hipGraphNodeType type_;
   std::vector<amd::Command*> commands_;
+  std::vector<Node> edges_;
+  std::vector<Node> dependencies_;
   bool visited_;
+  // count of in coming edges
+  size_t inDegree_;
+  // count of outgoing edges
+  size_t outDegree_;
+  static int nextID;
+  struct ihipGraph* parentGraph_;
 
  public:
   hipGraphNode(hipGraphNodeType type) {
     type_ = type;
     level_ = 0;
     visited_ = false;
+    inDegree_ = 0;
+    outDegree_ = 0;
+    id_ = nextID++;
+    parentGraph_ = nullptr;
+  }
+  /// Copy Constructor
+  hipGraphNode(const hipGraphNode& node) {
+    level_ = node.level_;
+    type_ = node.type_;
+    inDegree_ = node.inDegree_;
+    outDegree_ = node.outDegree_;
+    visited_ = false;
+    id_ = node.id_;
+    parentGraph_ = nullptr;
   }
   virtual ~hipGraphNode() {
-    for (auto command : commands_) {
-      delete command;
+    for (auto node : edges_) {
+      node->RemoveDependency(this);
+    }
+    for (auto node : dependencies_) {
+      node->RemoveEdge(this);
     }
   }
+  /// Create amd::command for the graph node
   virtual hipError_t CreateCommand(amd::HostQueue* queue) { return hipSuccess; }
+  /// Method to release amd::command part of node
+  virtual void ReleaseCommand() {
+    for (auto command : commands_) {
+      command->release();
+    }
+    commands_.clear();
+  }
+  /// Return node unique ID
+  int GetID() const { return id_; }
+  /// Returns command for graph node
   std::vector<amd::Command*>& GetCommands() { return commands_; }
-  hipGraphNodeType GetType() { return type_; }
-  uint32_t GetLevel() { return level_; }
+  /// Returns graph node type
+  hipGraphNodeType GetType() const { return type_; }
+  /// Returns graph node in coming edges
+  uint32_t GetLevel() const { return level_; }
+  /// Set graph node level
   void SetLevel(uint32_t level) { level_ = level; }
+  /// Clone graph node
+  virtual hipGraphNode* clone() const = 0;
+  /// Returns graph node indegree
+  size_t GetInDegree() const { return inDegree_; }
+  /// Updates indegree of the node
+  void SetInDegree(size_t inDegree) { inDegree_ = inDegree; }
+  /// Returns graph node outdegree
+  size_t GetOutDegree() const { return outDegree_; }
+  ///  Updates outdegree of the node
+  void SetOutDegree(size_t outDegree) { outDegree_ = outDegree; }
+  /// Returns graph node dependencies
+  const std::vector<Node>& GetDependencies() const { return dependencies_; }
+  /// Update graph node dependecies
+  void SetDependencies(std::vector<Node>& dependencies) {
+    for (auto entry : dependencies) {
+      dependencies_.push_back(entry);
+    }
+  }
+  /// Add graph node dependency
+  void AddDependency(const Node& node) { dependencies_.push_back(node); }
+  /// Remove graph node dependency
+  void RemoveDependency(const Node& node) {
+    dependencies_.erase(std::remove(dependencies_.begin(), dependencies_.end(), node),
+                        dependencies_.end());
+  }
+  /// Return graph node children
+  const std::vector<Node>& GetEdges() const { return edges_; }
+  /// Updates graph node children
+  void SetEdges(std::vector<Node>& edges) {
+    for (auto entry : edges) {
+      edges_.push_back(entry);
+    }
+  }
+  /// Add edge, update parent node outdegree, child node indegree, level and dependency
+  void AddEdge(const Node& childNode) {
+    edges_.push_back(childNode);
+    outDegree_++;
+    childNode->SetInDegree(childNode->GetInDegree() + 1);
+    childNode->SetLevel(std::max(childNode->GetLevel(), GetLevel() + 1));
+    childNode->AddDependency(this);
+  }
+  /// Remove edge, update parent node outdegree, child node indegree, level and dependency
+  void RemoveEdge(const Node& childNode) {
+    edges_.erase(std::remove(edges_.begin(), edges_.end(), childNode), edges_.end());
+    outDegree_--;
+    childNode->SetInDegree(childNode->GetInDegree() - 1);
+    const std::vector<Node>& dependencies = childNode->GetDependencies();
+    int32_t level = 0;
+    int32_t parentLevel = 0;
+    for (auto parent : dependencies) {
+      parentLevel = parent->GetLevel();
+      level = std::max(level, (parentLevel + 1));
+    }
+    childNode->SetLevel(level);
+    childNode->RemoveDependency(this);
+  }
+  /// Get Runlist of the nodes embedded as part of the graphnode(e.g. ChildGraph)
+  virtual void GetRunList(std::vector<std::vector<Node>>& parallelList,
+                          std::unordered_map<Node, std::vector<Node>>& dependencies) {}
+  /// Get levelorder of the nodes embedded as part of the graphnode(e.g. ChildGraph)
+  virtual void LevelOrder(std::vector<Node>& levelOrder) {}
+  /// Update waitlist of the nodes embedded as part of the graphnode(e.g. ChildGraph)
+  virtual void UpdateEventWaitLists() {}
+  /// Enqueue commands part of the node
+  virtual void EnqueueCommands() {
+    for (auto& command : commands_) {
+      command->enqueue();
+    }
+  }
+  /// Reset commands part of the node
+  virtual void ResetStatus() {
+    for (auto& command : commands_) {
+      command->resetStatus(CL_INT_MAX);
+    }
+  }
+  ihipGraph* GetParentGraph() { return parentGraph_; }
+  void SetParentGraph(ihipGraph* graph) { parentGraph_ = graph; }
 };
 
 struct ihipGraph {
-  std::unordered_map<Node, size_t> nodeInDegree_;   // count of in coming edges for every vertex
-  std::unordered_map<Node, size_t> nodeOutDegree_;  // count of outgoing edges for every vertex
   std::vector<Node> vertices_;
-  std::unordered_map<Node, std::vector<Node>> edges_;
 
  public:
   ihipGraph() {}
-  ~ihipGraph(){};
+  ~ihipGraph() {
+    for (auto node : vertices_) {
+      delete node;
+    }
+  }
   /// add node to the graph
-  hipError_t AddNode(const Node& node);
-  /// add edge to the graph
-  hipError_t AddEdge(const Node& parentNode, const Node& childNode);
+  void AddNode(const Node& node);
+  void RemoveNode(const Node& node);
   /// Returns root nodes, all vertices with 0 in-degrees
   std::vector<Node> GetRootNodes() const;
   /// Returns leaf nodes, all vertices with 0 out-degrees
@@ -81,7 +197,7 @@ struct ihipGraph {
   /// Returns total numbers of nodes in the graph
   size_t GetNodeCount() const { return vertices_.size(); }
   /// returns all the nodes in the graph
-  std::vector<Node> GetNodes() const { return vertices_; }
+  const std::vector<Node>& GetNodes() const { return vertices_; }
   /// returns all the edges in the graph
   std::vector<std::pair<Node, Node>> GetEdges() const;
   void GetRunListUtil(Node v, std::unordered_map<Node, bool>& visited,
@@ -89,7 +205,101 @@ struct ihipGraph {
                       std::unordered_map<Node, std::vector<Node>>& dependencies);
   void GetRunList(std::vector<std::vector<Node>>& parallelList,
                   std::unordered_map<Node, std::vector<Node>>& dependencies);
-  hipError_t LevelOrder(std::vector<Node>& levelOrder);
+  void LevelOrder(std::vector<Node>& levelOrder);
+  ihipGraph* clone() const;
+};
+
+struct hipChildGraphNode : public hipGraphNode {
+  struct ihipGraph* childGraph_;
+  std::vector<Node> childGraphlevelOrder_;
+
+ public:
+  hipChildGraphNode(ihipGraph* g) : hipGraphNode(hipGraphNodeTypeGraph) {
+    // ToDo: clone the child graph
+    childGraph_ = g->clone();
+  }
+
+  ~hipChildGraphNode() { delete childGraph_; }
+
+  hipChildGraphNode(const hipChildGraphNode& rhs) : hipGraphNode(rhs) {
+    childGraph_ = rhs.childGraph_->clone();
+  }
+
+  hipGraphNode* clone() const {
+    return new hipChildGraphNode(static_cast<hipChildGraphNode const&>(*this));
+  }
+
+  ihipGraph* GetChildGraph() { return childGraph_; }
+
+  hipError_t CreateCommand(amd::HostQueue* queue) {
+    commands_.reserve(2);
+    amd::Command::EventWaitList eventWaitList;
+    // Command for start of the graph
+    commands_.push_back(new amd::Marker(*queue, false, eventWaitList));
+    // Command for end of the graph
+    commands_.push_back(new amd::Marker(*queue, false, eventWaitList));
+    childGraph_->LevelOrder(childGraphlevelOrder_);
+    return hipSuccess;
+  }
+
+  void UpdateEventWaitLists() {
+    // ChildGraph should start after all parents
+    if (commands_.size() == 2) {
+      std::vector<Node> rootNodes = childGraph_->GetRootNodes();
+      amd::Command::EventWaitList waitList;
+      waitList.push_back(commands_[0]);
+      for (auto& node : rootNodes) {
+        for (auto command : node->GetCommands()) {
+          command->updateEventWaitList(waitList);
+        }
+      }
+      waitList.clear();
+      // End command should wait for graph to finish
+      std::vector<Node> leafNodes = childGraph_->GetLeafNodes();
+      for (auto& node : leafNodes) {
+        for (auto command : node->GetCommands()) {
+          waitList.push_back(command);
+        }
+      }
+      commands_[1]->updateEventWaitList(waitList);
+    } else {
+      ClPrint(amd::LOG_ERROR, amd::LOG_CODE,
+              "[hipGraph] childgraph node commands are not created!\n");
+    }
+  }
+
+  void ResetStatus() {
+    if (commands_.size() == 2) {
+      commands_[0]->resetStatus(CL_INT_MAX);
+      commands_[1]->resetStatus(CL_INT_MAX);
+    }
+    for (auto& node : childGraphlevelOrder_) {
+      node->ResetStatus();
+    }
+  }
+
+  void GetRunList(std::vector<std::vector<Node>>& parallelList,
+                  std::unordered_map<Node, std::vector<Node>>& dependencies) {
+    childGraph_->GetRunList(parallelList, dependencies);
+  }
+
+  void LevelOrder(std::vector<Node>& levelOrder) { childGraph_->LevelOrder(levelOrder); }
+
+  void EnqueueCommands() {
+    if (commands_.size() == 2) {
+      // enqueue child graph start command
+      commands_[0]->enqueue();
+      // enqueue nodes in child graph in level order
+      for (auto& node : childGraphlevelOrder_) {
+        node->EnqueueCommands();
+      }
+      // enqueue child graph end command
+      commands_[1]->enqueue();
+    } else {
+      ClPrint(amd::LOG_ERROR, amd::LOG_CODE,
+              "[hipGraph] childgraph node commands are not created!\n");
+    }
+  }
 };
 
 class hipGraphKernelNode : public hipGraphNode {
@@ -103,7 +313,13 @@ class hipGraphKernelNode : public hipGraphNode {
     func_ = func;
   }
   ~hipGraphKernelNode() { delete pKernelParams_; }
-
+  hipGraphKernelNode(const hipGraphKernelNode& rhs) : hipGraphNode(rhs) {
+    pKernelParams_ = new hipKernelNodeParams(*rhs.pKernelParams_);
+    func_ = rhs.func_;
+  }
+  hipGraphNode* clone() const {
+    return new hipGraphKernelNode(static_cast<hipGraphKernelNode const&>(*this));
+  }
   hipError_t CreateCommand(amd::HostQueue* queue) {
     commands_.reserve(1);
     amd::Command* command;
@@ -153,6 +369,14 @@ class hipGraphMemcpyNode : public hipGraphNode {
   }
   ~hipGraphMemcpyNode() { delete pCopyParams_; }
 
+  hipGraphMemcpyNode(const hipGraphMemcpyNode& rhs) : hipGraphNode(rhs) {
+    pCopyParams_ = new hipMemcpy3DParms(*rhs.pCopyParams_);
+  }
+
+  hipGraphNode* clone() const {
+    return new hipGraphMemcpyNode(static_cast<hipGraphMemcpyNode const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue) {
     commands_.reserve(1);
     amd::Command* command;
@@ -178,8 +402,13 @@ class hipGraphMemcpyNode1D : public hipGraphNode {
  public:
   hipGraphMemcpyNode1D(void* dst, const void* src, size_t count, hipMemcpyKind kind)
       : hipGraphNode(hipGraphNodeTypeMemcpy1D), dst_(dst), src_(src), count_(count), kind_(kind) {}
+
   ~hipGraphMemcpyNode1D() {}
 
+  hipGraphNode* clone() const {
+    return new hipGraphMemcpyNode1D(static_cast<hipGraphMemcpyNode1D const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue) {
     commands_.reserve(1);
     amd::Command* command = nullptr;
@@ -214,6 +443,11 @@ template <class T> class hipGraphMemcpyNodeFromSymbol : public hipGraphNode {
         kind_(kind) {}
   ~hipGraphMemcpyNodeFromSymbol() {}
 
+  hipGraphNode* clone() const {
+    return new hipGraphMemcpyNodeFromSymbol(
+        static_cast<hipGraphMemcpyNodeFromSymbol const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue);
 
   void SetParams(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind) {
@@ -243,6 +477,10 @@ template <class T> class hipGraphMemcpyNodeToSymbol : public hipGraphNode {
         kind_(kind) {}
   ~hipGraphMemcpyNodeToSymbol() {}
 
+  hipGraphNode* clone() const {
+    return new hipGraphMemcpyNodeToSymbol(static_cast<hipGraphMemcpyNodeToSymbol const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue);
 
   void SetParams(const T& symbol, void* src, size_t count, size_t offset, hipMemcpyKind kind) {
@@ -262,6 +500,14 @@ class hipGraphMemsetNode : public hipGraphNode {
     pMemsetParams_ = new hipMemsetParams(*pMemsetParams);
   }
   ~hipGraphMemsetNode() { delete pMemsetParams_; }
+  // Copy constructor
+  hipGraphMemsetNode(const hipGraphMemsetNode& memsetNode) : hipGraphNode(memsetNode) {
+    pMemsetParams_ = new hipMemsetParams(*memsetNode.pMemsetParams_);
+  }
+
+  hipGraphNode* clone() const {
+    return new hipGraphMemsetNode(static_cast<hipGraphMemsetNode const&>(*this));
+  }
 
   hipError_t CreateCommand(amd::HostQueue* queue) {
     if (pMemsetParams_->height == 1) {
@@ -294,6 +540,10 @@ class hipGraphEventRecordNode : public hipGraphNode {
       : hipGraphNode(hipGraphNodeTypeEventRecord), event_(event) {}
   ~hipGraphEventRecordNode() {}
 
+  hipGraphNode* clone() const {
+    return new hipGraphEventRecordNode(static_cast<hipGraphEventRecordNode const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue);
 
   void GetParams(hipEvent_t* event) { *event = event_; }
@@ -308,6 +558,10 @@ class hipGraphEventWaitNode : public hipGraphNode {
       : hipGraphNode(hipGraphNodeTypeWaitEvent), event_(event) {}
   ~hipGraphEventWaitNode() {}
 
+  hipGraphNode* clone() const {
+    return new hipGraphEventWaitNode(static_cast<hipGraphEventWaitNode const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue);
 
   void GetParams(hipEvent_t* event) { *event = event_; }
@@ -323,6 +577,14 @@ class hipGraphHostNode : public hipGraphNode {
   }
   ~hipGraphHostNode() { delete pNodeParams_; }
 
+  hipGraphHostNode(const hipGraphHostNode& hostNode) : hipGraphNode(hostNode) {
+    pNodeParams_ = new hipHostNodeParams(*hostNode.pNodeParams_);
+  }
+
+  hipGraphNode* clone() const {
+    return new hipGraphHostNode(static_cast<hipGraphHostNode const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue);
 
   void GetParams(hipHostNodeParams* params) {
@@ -338,6 +600,10 @@ class hipGraphEmptyNode : public hipGraphNode {
   hipGraphEmptyNode() : hipGraphNode(hipGraphNodeTypeEmpty) {}
   ~hipGraphEmptyNode() {}
 
+  hipGraphNode* clone() const {
+    return new hipGraphEmptyNode(static_cast<hipGraphEmptyNode const&>(*this));
+  }
+
   hipError_t CreateCommand(amd::HostQueue* queue) {
     amd::Command::EventWaitList waitList;
     commands_.reserve(1);
@@ -349,6 +615,7 @@ class hipGraphEmptyNode : public hipGraphNode {
 
 struct hipGraphExec {
   std::vector<std::vector<Node>> parallelLists_;
+  // level order of the graph doesn't include nodes embedded as part of the child graph
   std::vector<Node> levelOrder_;
   std::unordered_map<Node, std::vector<Node>> nodeWaitLists_;
   std::vector<amd::HostQueue*> parallelQueues_;
@@ -374,14 +641,14 @@ struct hipGraphExec {
       queue->release();
     }
     for (auto node : levelOrder_) {
-      delete node;
+      node->ReleaseCommand();
     }
   }
 
   hipError_t CreateQueues();
   hipError_t FillCommands();
   hipError_t Init();
-  hipError_t UpdateGraphToWaitOnRoot();
+  void UpdateGraphToWaitOnRoot();
   hipError_t Run(hipStream_t stream);
   static void ResetGraph(cl_event event, cl_int command_exec_status, void* user_data);
 };

projects/clr/hipamd/src/hip_internal.hpp

@@ -177,11 +177,13 @@ namespace hip {
     hipStream_t parentStream_;
     /// Last graph node captured in the stream
     std::vector<hipGraphNode_t> lastCapturedNodes_;
+    /// dependencies removed via API hipStreamUpdateCaptureDependencies
+    std::vector<hipGraphNode_t> removedDependencies_;
     /// Derived streams/Paralell branches from the origin stream
     std::vector<hipStream_t> parallelCaptureStreams_;
     /// Capture events
     std::vector<hipEvent_t> captureEvents_;
-
+    unsigned long long captureID_;
   public:
     Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false,
            const std::vector<uint32_t>& cuMask = {},
@@ -231,8 +233,19 @@ namespace hip {
       lastCapturedNodes_.clear();
       lastCapturedNodes_.push_back(graphNode);
     }
+    /// returns updated dependencies removed
+    const std::vector<hipGraphNode_t>& GetRemovedDependencies() {
+      return removedDependencies_;
+    }
     /// Append captured node via the wait event cross stream
-    void AddCrossCapturedNode(std::vector<hipGraphNode_t> graphNodes) {
+    void AddCrossCapturedNode(std::vector<hipGraphNode_t> graphNodes, bool replace = false) {
+      // replace dependencies as per flag hipStreamSetCaptureDependencies
+      if (replace == true) {
+        for (auto node : lastCapturedNodes_) {
+          removedDependencies_.push_back(node);
+        }
+        lastCapturedNodes_.clear();
+      }
       for (auto node : graphNodes) {
         lastCapturedNodes_.push_back(node);
       }
@@ -241,6 +254,8 @@ namespace hip {
     void SetCaptureGraph(hipGraph_t pGraph) {
       pCaptureGraph_ = pGraph;
       captureStatus_ = hipStreamCaptureStatusActive;
+      // ID is generated in Begin Capture i.e.. when capture status is active
+      captureID_ = GenerateCaptureID();
     }
     /// reset capture parameters
     hipError_t EndCapture();
@@ -252,6 +267,13 @@ namespace hip {
     void SetParentStream(hipStream_t parentStream) { parentStream_ = parentStream; }
     /// Get parent stream
     hipStream_t GetParentStream() { return parentStream_; }
+    /// Generate ID for stream capture unique over the lifetime of the process
+    static int GenerateCaptureID() {
+      static std::atomic<unsigned long long> uid(0);
+      return ++uid;
+    }
+    /// Get Capture ID
+    int GetCaptureID() { return captureID_; }
   };
 
   /// HIP Device class

projects/clr/hipamd/src/hip_stream.cpp

@@ -41,7 +41,8 @@ Stream::Stream(hip::Device* dev, Priority p, unsigned int f, bool null_stream,
       null_(null_stream),
       cuMask_(cuMask),
       captureStatus_(captureStatus),
-      originStream_(false) {}
+      originStream_(false),
+      captureID_(0) {}
 
 // ================================================================================================
 Stream::~Stream() {