SWDEV-301667 - Port optimization to save extra packet to graphs

Change-Id: Ibaf64a4efe070c42620e6e153c1862a4a0b15664


[ROCm/clr commit: e1e5d071ba]

projects/clr/hipamd/src/hip_event.cpp

@@ -383,8 +383,6 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
 }
 
 hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node EventRecord on stream : %p, Event %p", stream, event);
   hipError_t status = hipSuccess;
   if (event == nullptr) {
     return hipErrorInvalidHandle;
@@ -398,6 +396,8 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream) {
   hip::Stream* hip_stream = hip::getStream(stream);
   e->SetCaptureStream(stream);
   if ((s != nullptr) && (s->GetCaptureStatus() == hipStreamCaptureStatusActive)) {
+    ClPrint(amd::LOG_INFO, amd::LOG_API,
+        "[hipGraph] Current capture node EventRecord on stream : %p, Event %p", stream, event);
     s->SetCaptureEvent(event);
     std::vector<hip::GraphNode*> lastCapturedNodes = s->GetLastCapturedNodes();
     if (!lastCapturedNodes.empty()) {

projects/clr/hipamd/src/hip_graph.cpp

@@ -63,7 +63,9 @@ inline hipError_t ihipGraphAddNode(hip::GraphNode* graphNode, hip::Graph* graph,
 
 hipError_t ihipGraphAddKernelNode(hip::GraphNode** pGraphNode, hip::Graph* graph,
                                   hip::GraphNode* const* pDependencies, size_t numDependencies,
-                                  const hipKernelNodeParams* pNodeParams, bool capture = true) {
+                                  const hipKernelNodeParams* pNodeParams,
+                                  const ihipExtKernelEvents* pNodeEvents = nullptr,
+                                  bool capture = true) {
   if (pGraphNode == nullptr || graph == nullptr ||
       (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr ||
       pNodeParams->func == nullptr) {
@@ -94,7 +96,7 @@ hipError_t ihipGraphAddKernelNode(hip::GraphNode** pGraphNode, hip::Graph* graph
     return hipErrorInvalidConfiguration;
   }
 
-  *pGraphNode = new hip::GraphKernelNode(pNodeParams);
+  *pGraphNode = new hip::GraphKernelNode(pNodeParams, pNodeEvents);
   status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture);
   return status;
 }
@@ -175,7 +177,7 @@ hipError_t ihipGraphAddMemsetNode(hip::GraphNode** pGraphNode, hip::Graph* graph
 hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction, dim3& gridDim,
                                   dim3& blockDim, void**& args, size_t& sharedMemBytes) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node kernel launch on stream : %p", stream);
+          "[hipGraph] Current capture node LaunchKernel on stream : %p", stream);
 
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
@@ -212,7 +214,15 @@ hipError_t ihipExtLaunchKernel(hipStream_t stream, hipFunction_t f, uint32_t glo
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
 
   hip::GraphNode* pGraphNode;
-  hipError_t status;
+  hipError_t status = hipSuccess;
+
+  // Consider start event as an explicit node because it is recorded as an event
+  // For the purpose of optimization, the stop event isn't a separate node as we bind it later
+  // to the KernelLaunch command. Pass only the stopEvent further.
+  ihipExtKernelEvents nodeEvents;
+  nodeEvents.startEvent_ = nullptr;
+  nodeEvents.stopEvent_ = stopEvent;
+
   if (startEvent != nullptr) {
     pGraphNode = new hip::GraphEventRecordNode(startEvent);
     status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(),
@@ -222,6 +232,7 @@ hipError_t ihipExtLaunchKernel(hipStream_t stream, hipFunction_t f, uint32_t glo
     }
     s->SetLastCapturedNode(pGraphNode);
   }
+
   hipKernelNodeParams nodeParams;
   nodeParams.func = f;
   nodeParams.blockDim = dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ);
@@ -230,23 +241,16 @@ hipError_t ihipExtLaunchKernel(hipStream_t stream, hipFunction_t f, uint32_t glo
                             globalWorkSizeZ / localWorkSizeZ);
   nodeParams.kernelParams = kernelParams;
   nodeParams.sharedMemBytes = sharedMemBytes;
+
   status =
       ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(),
-                             s->GetLastCapturedNodes().size(), &nodeParams);
+                             s->GetLastCapturedNodes().size(), &nodeParams, &nodeEvents);
 
   if (status != hipSuccess) {
     return status;
   }
   s->SetLastCapturedNode(pGraphNode);
-  if (stopEvent != nullptr) {
-    pGraphNode = new hip::GraphEventRecordNode(stopEvent);
-    status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(),
-                              s->GetLastCapturedNodes().size());
-    if (status != hipSuccess) {
-      return status;
-    }
-    s->SetLastCapturedNode(pGraphNode);
-  }
+
   return hipSuccess;
 }
 
@@ -258,7 +262,7 @@ hipError_t capturehipExtModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f
                                            void**& extra, hipEvent_t& startEvent,
                                            hipEvent_t& stopEvent, uint32_t& flags) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node Ext Module launch kernel on stream : %p", stream);
+          "[hipGraph] Current capture node ExtModuleLaunchKernel on stream : %p", stream);
   return ihipExtLaunchKernel(stream, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
                              localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes,
                              kernelParams, extra, startEvent, stopEvent, flags);
@@ -268,7 +272,7 @@ hipError_t capturehipExtLaunchKernel(hipStream_t& stream, const void*& hostFunct
                                      dim3& blockDim, void**& args, size_t& sharedMemBytes,
                                      hipEvent_t& startEvent, hipEvent_t& stopEvent, int& flags) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node Ext kernel launch on stream : %p", stream);
+          "[hipGraph] Current capture node ExtLaunchKernel on stream : %p", stream);
   return ihipExtLaunchKernel(
       stream, reinterpret_cast<hipFunction_t>(const_cast<void*>(hostFunction)),
       gridDim.x * blockDim.x, gridDim.y * blockDim.y, gridDim.z * blockDim.z, blockDim.x,
@@ -281,7 +285,7 @@ hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, u
                                         uint32_t& sharedMemBytes, void**& kernelParams,
                                         void**& extra) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node module launch kernel launch on stream : %p", stream);
+          "[hipGraph] Current capture node ModuleLaunchKernel on stream : %p", stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -306,7 +310,7 @@ hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, u
 }
 
 hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*& p) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy3D on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node Memcpy3D on stream : %p",
           stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
@@ -326,7 +330,7 @@ hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*&
 hipError_t capturehipMemcpy2DAsync(hipStream_t& stream, void*& dst, size_t& dpitch,
                                    const void*& src, size_t& spitch, size_t& width, size_t& height,
                                    hipMemcpyKind& kind) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy2D on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node Memcpy2D on stream : %p",
           stream);
   if (dst == nullptr || src == nullptr) {
     return hipErrorInvalidValue;
@@ -365,7 +369,7 @@ hipError_t capturehipMemcpy2DFromArrayAsync(hipStream_t& stream, void*& dst, siz
                                             size_t& hOffsetSrc, size_t& width, size_t& height,
                                             hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node Memcpy2DFromArray on stream : %p", stream);
+          "[hipGraph] Current capture node Memcpy2DFromArray on stream : %p", stream);
   if (src == nullptr || dst == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -400,7 +404,7 @@ hipError_t capturehipMemcpyFromArrayAsync(hipStream_t& stream, void*& dst, hipAr
                                           size_t& wOffsetSrc, size_t& hOffsetSrc, size_t& count,
                                           hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node Memcpy2DFromArray on stream : %p", stream);
+          "[hipGraph] Current capture node Memcpy2DFromArray on stream : %p", stream);
   if (src == nullptr || dst == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -438,7 +442,7 @@ hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray*& dst, s
                                           size_t& hOffset, const void*& src, size_t& spitch,
                                           size_t& width, size_t& height, hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node Memcpy2DFromArray on stream : %p", stream);
+          "[hipGraph] Current capture node Memcpy2DFromArray on stream : %p", stream);
   if (src == nullptr || dst == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -473,7 +477,7 @@ hipError_t capturehipMemcpyToArrayAsync(hipStream_t& stream, hipArray_t& dst, si
                                         size_t& hOffset, const void*& src, size_t& count,
                                         hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node Memcpy2DFromArray on stream : %p", stream);
+          "[hipGraph] Current capture node Memcpy2DFromArray on stream : %p", stream);
   if (src == nullptr || dst == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -509,7 +513,7 @@ hipError_t capturehipMemcpyToArrayAsync(hipStream_t& stream, hipArray_t& dst, si
 
 hipError_t capturehipMemcpyParam2DAsync(hipStream_t& stream, const hip_Memcpy2D*& pCopy) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node MemcpyParam2D on stream : %p", stream);
+          "[hipGraph] Current capture node MemcpyParam2D on stream : %p", stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -558,7 +562,7 @@ hipError_t capturehipMemcpyParam2DAsync(hipStream_t& stream, const hip_Memcpy2D*
 hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArray*& srcArray,
                                      size_t& srcOffset, size_t& ByteCount) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node MemcpyParam2D on stream : %p", stream);
+          "[hipGraph] Current capture node MemcpyParam2D on stream : %p", stream);
   if (srcArray == nullptr || dstHost == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -586,7 +590,7 @@ hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArr
 hipError_t capturehipMemcpyHtoAAsync(hipStream_t& stream, hipArray*& dstArray, size_t& dstOffset,
                                      const void*& srcHost, size_t& ByteCount) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node MemcpyParam2D on stream : %p", stream);
+          "[hipGraph] Current capture node MemcpyParam2D on stream : %p", stream);
   if (dstArray == nullptr || srcHost == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -635,7 +639,7 @@ hipError_t capturehipMemcpy(hipStream_t stream, void* dst, const void* src, size
 
 hipError_t capturehipMemcpyAsync(hipStream_t& stream, void*& dst, const void*& src,
                                  size_t& sizeBytes, hipMemcpyKind& kind) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy1D on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node Memcpy1D on stream : %p",
           stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
@@ -645,7 +649,7 @@ hipError_t capturehipMemcpyAsync(hipStream_t& stream, void*& dst, const void*& s
 
 hipError_t capturehipMemcpyHtoDAsync(hipStream_t& stream, hipDeviceptr_t& dstDevice, void*& srcHost,
                                      size_t& ByteCount, hipMemcpyKind& kind) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node MemcpyHtoD on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node MemcpyHtoD on stream : %p",
           stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
@@ -657,7 +661,7 @@ hipError_t capturehipMemcpyDtoDAsync(hipStream_t& stream, hipDeviceptr_t& dstDev
                                      hipDeviceptr_t& srcDevice, size_t& ByteCount,
                                      hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node hipMemcpyDtoD on stream : %p", stream);
+          "[hipGraph] Current capture node hipMemcpyDtoD on stream : %p", stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -667,7 +671,7 @@ hipError_t capturehipMemcpyDtoDAsync(hipStream_t& stream, hipDeviceptr_t& dstDev
 hipError_t capturehipMemcpyDtoHAsync(hipStream_t& stream, void*& dstHost, hipDeviceptr_t& srcDevice,
                                      size_t& ByteCount, hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node hipMemcpyDtoH on stream : %p", stream);
+          "[hipGraph] Current capture node hipMemcpyDtoH on stream : %p", stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -677,7 +681,7 @@ hipError_t capturehipMemcpyDtoHAsync(hipStream_t& stream, void*& dstHost, hipDev
 hipError_t capturehipMemcpyFromSymbolAsync(hipStream_t& stream, void*& dst, const void*& symbol,
                                            size_t& sizeBytes, size_t& offset, hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node MemcpyFromSymbolNode on stream : %p", stream);
+          "[hipGraph] Current capture node MemcpyFromSymbolNode on stream : %p", stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -703,7 +707,7 @@ hipError_t capturehipMemcpyFromSymbolAsync(hipStream_t& stream, void*& dst, cons
 hipError_t capturehipMemcpyToSymbolAsync(hipStream_t& stream, const void*& symbol, const void*& src,
                                          size_t& sizeBytes, size_t& offset, hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node MemcpyToSymbolNode on stream : %p", stream);
+          "[hipGraph] Current capture node MemcpyToSymbolNode on stream : %p", stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -726,7 +730,7 @@ hipError_t capturehipMemcpyToSymbolAsync(hipStream_t& stream, const void*& symbo
 
 hipError_t capturehipMemsetAsync(hipStream_t& stream, void*& dst, int& value, size_t& valueSize,
                                  size_t& sizeBytes) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset1D on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node Memset1D on stream : %p",
           stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
@@ -752,7 +756,7 @@ hipError_t capturehipMemsetAsync(hipStream_t& stream, void*& dst, int& value, si
 
 hipError_t capturehipMemset2DAsync(hipStream_t& stream, void*& dst, size_t& pitch, int& value,
                                    size_t& width, size_t& height) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset2D on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node Memset2D on stream : %p",
           stream);
   hipMemsetParams memsetParams = {0};
   if (!hip::isValid(stream)) {
@@ -777,7 +781,7 @@ hipError_t capturehipMemset2DAsync(hipStream_t& stream, void*& dst, size_t& pitc
 
 hipError_t capturehipMemset3DAsync(hipStream_t& stream, hipPitchedPtr& pitchedDevPtr, int& value,
                                    hipExtent& extent) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset3D on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node Memset3D on stream : %p",
           stream);
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
@@ -786,7 +790,7 @@ hipError_t capturehipMemset3DAsync(hipStream_t& stream, hipPitchedPtr& pitchedDe
 }
 
 hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*& userData) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node host on stream : %p",
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node host on stream : %p",
           stream);
   if (fn == nullptr) {
     return hipErrorInvalidValue;
@@ -1090,7 +1094,8 @@ hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
   hip::GraphNode* node;
   hipError_t status = ihipGraphAddKernelNode(
       &node, reinterpret_cast<hip::Graph*>(graph),
-      reinterpret_cast<hip::GraphNode* const*>(pDependencies), numDependencies, pNodeParams, false);
+      reinterpret_cast<hip::GraphNode* const*>(pDependencies), numDependencies, pNodeParams,
+      nullptr, false);
   *pGraphNode = reinterpret_cast<hipGraphNode*>(node);
   HIP_RETURN(status);
 }
@@ -1107,7 +1112,8 @@ hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
   hip::GraphNode* node;
   hipError_t status = ihipGraphAddMemcpyNode(
       &node, reinterpret_cast<hip::Graph*>(graph),
-      reinterpret_cast<hip::GraphNode* const*>(pDependencies), numDependencies, pCopyParams, false);
+      reinterpret_cast<hip::GraphNode* const*>(pDependencies), numDependencies, pCopyParams,
+      false);
   *pGraphNode = reinterpret_cast<hipGraphNode_t>(node);
   HIP_RETURN(status);
 }
@@ -1115,16 +1121,16 @@ hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
 hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                    const hipGraphNode_t* pDependencies, size_t numDependencies,
                                    void* dst, const void* src, size_t count, hipMemcpyKind kind) {
-  HIP_INIT_API(hipGraphAddMemcpyNode1D, pGraphNode, graph, pDependencies, numDependencies, dst, src,
-               count, kind);
+  HIP_INIT_API(hipGraphAddMemcpyNode1D, pGraphNode, graph, pDependencies, numDependencies, dst,
+               src, count, kind);
   if (pGraphNode == nullptr || graph == nullptr ||
       (numDependencies > 0 && pDependencies == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   hip::GraphNode* node;
   hipError_t status = ihipGraphAddMemcpyNode1D(&node, reinterpret_cast<hip::Graph*>(graph),
-                                               reinterpret_cast<hip::GraphNode* const*>(pDependencies),
-                                               numDependencies, dst, src, count, kind, false);
+                        reinterpret_cast<hip::GraphNode* const*>(pDependencies),
+                        numDependencies, dst, src, count, kind, false);
   *pGraphNode = reinterpret_cast<hipGraphNode_t>(node);
   HIP_RETURN(status);
 }
@@ -1155,7 +1161,8 @@ hipError_t hipGraphExecMemcpyNodeSetParams1D(hipGraphExec_t hGraphExec, hipGraph
   if (clonedNode == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  HIP_RETURN(reinterpret_cast<hip::GraphMemcpyNode1D*>(clonedNode)->SetParams(dst, src, count, kind));
+  HIP_RETURN(reinterpret_cast<hip::GraphMemcpyNode1D*>(clonedNode)->SetParams(dst, src,
+                                                                              count, kind));
 }
 
 hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
@@ -1172,7 +1179,7 @@ hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
       ihipGraphAddMemsetNode(&node, reinterpret_cast<hip::Graph*>(graph),
                              reinterpret_cast<hip::GraphNode* const*>(pDependencies),
                              numDependencies, pMemsetParams, false);
-  *pGraphNode = reinterpret_cast<hipGraphNode_t>(node);                             
+  *pGraphNode = reinterpret_cast<hipGraphNode_t>(node);
   HIP_RETURN(status);
 }
 
@@ -1233,7 +1240,8 @@ hipError_t ihipGraphInstantiate(hip::GraphExec** pGraphExec, hip::Graph* graph,
     return hipErrorInvalidValue;
   }
   *pGraphExec =
-      new hip::GraphExec(graphNodes, parallelLists, nodeWaitLists, clonedGraph, clonedNodes, flags);
+      new hip::GraphExec(graphNodes, parallelLists, nodeWaitLists, clonedGraph, clonedNodes,
+                         flags);
   if (*pGraphExec != nullptr) {
     graph->SetGraphInstantiated(true);
     return (*pGraphExec)->Init();
@@ -1378,7 +1386,8 @@ hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams*
 hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node,
                                        const hipKernelNodeParams* pNodeParams) {
   HIP_INIT_API(hipGraphKernelNodeSetParams, node, pNodeParams);
-  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) || pNodeParams == nullptr || pNodeParams->func == nullptr) {
+  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) ||
+                                   pNodeParams == nullptr || pNodeParams->func == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   HIP_RETURN(reinterpret_cast<hip::GraphKernelNode*>(node)->SetParams(pNodeParams));
@@ -1386,7 +1395,8 @@ hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node,
 
 hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams) {
   HIP_INIT_API(hipGraphMemcpyNodeGetParams, node, pNodeParams);
-  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) || pNodeParams == nullptr) {
+  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) ||
+                                   pNodeParams == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   reinterpret_cast<hip::GraphMemcpyNode*>(node)->GetParams(pNodeParams);
@@ -1421,7 +1431,8 @@ hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt
 
 hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams) {
   HIP_INIT_API(hipGraphMemcpyNodeSetParams, node, pNodeParams);
-  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) || pNodeParams == nullptr) {
+  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) ||
+                                   pNodeParams == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   HIP_RETURN(reinterpret_cast<hip::GraphMemcpyNode*>(node)->SetParams(pNodeParams));
@@ -1431,7 +1442,8 @@ hipError_t hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo
                                            hipMemcpy3DParms* pNodeParams) {
   HIP_INIT_API(hipGraphExecMemcpyNodeSetParams, hGraphExec, node, pNodeParams);
   hip::GraphNode* n = reinterpret_cast<hip::GraphNode*>(node);
-  if (hGraphExec == nullptr || !hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(n))) {
+  if (hGraphExec == nullptr ||
+                    !hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(n))) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   if (ihipMemcpy3D_validate(pNodeParams) != hipSuccess) {
@@ -1451,7 +1463,8 @@ hipError_t hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo
 
 hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams) {
   HIP_INIT_API(hipGraphMemsetNodeGetParams, node, pNodeParams);
-  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) || pNodeParams == nullptr) {
+  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) ||
+                                   pNodeParams == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   reinterpret_cast<hip::GraphMemsetNode*>(node)->GetParams(pNodeParams);
@@ -1460,7 +1473,8 @@ hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNo
 
 hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams) {
   HIP_INIT_API(hipGraphMemsetNodeSetParams, node, pNodeParams);
-  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) || pNodeParams == nullptr) {
+  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) ||
+                                   pNodeParams == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   if (pNodeParams->height > 1 &&
@@ -1751,7 +1765,8 @@ hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode
       (to == nullptr && from != nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  const std::vector<std::pair<hip::GraphNode*, hip::GraphNode*>> edges = reinterpret_cast<hip::Graph*>(graph)->GetEdges();
+  const std::vector<std::pair<hip::GraphNode*, hip::GraphNode*>> edges =
+                                                reinterpret_cast<hip::Graph*>(graph)->GetEdges();
   // returns only the number of edges in numEdges when from and to are null
   if (from == nullptr && to == nullptr) {
     *numEdges = edges.size();
@@ -1797,8 +1812,8 @@ hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDep
     for (int i = 0; i < dependencies.size(); i++) {
       pDependencies[i] = reinterpret_cast<hipGraphNode_t>(dependencies[i]);
     }
-    // pNumDependencies > actual number of dependencies, the remaining entries in pDependencies will
-    // be set to NULL
+    // 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;
     }
@@ -1975,7 +1990,8 @@ 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 ||
-      !hip::Graph::isGraphValid(reinterpret_cast<hip::Graph*>(graph)) || (pDependencies == nullptr && numDependencies > 0)) {
+      !hip::Graph::isGraphValid(reinterpret_cast<hip::Graph*>(graph)) ||
+                                (pDependencies == nullptr && numDependencies > 0)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   size_t sym_size = 0;
@@ -2022,7 +2038,8 @@ hipError_t hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hi
     HIP_RETURN(hipErrorInvalidSymbol);
   }
   hip::GraphNode* n = reinterpret_cast<hip::GraphNode*>(node);
-  if (hGraphExec == nullptr || src == nullptr || !hip::GraphNode::isNodeValid(n) || count == 0 || src == symbol) {
+  if (hGraphExec == nullptr || src == nullptr || !hip::GraphNode::isNodeValid(n) || count == 0 ||
+      src == symbol) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
@@ -2163,7 +2180,8 @@ hipError_t hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
 
 hipError_t hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNodeParams) {
   HIP_INIT_API(hipGraphHostNodeGetParams, node, pNodeParams);
-  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) || pNodeParams == nullptr) {
+  if (!hip::GraphNode::isNodeValid(reinterpret_cast<hip::GraphNode*>(node)) ||
+                                   pNodeParams == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   reinterpret_cast<hip::GraphHostNode*>(node)->GetParams(pNodeParams);

projects/clr/hipamd/src/hip_graph_internal.cpp

@@ -285,7 +285,7 @@ std::vector<Node> Graph::GetRootNodes() const {
   for (auto entry : vertices_) {
     if (entry->GetInDegree() == 0) {
       roots.push_back(entry);
-      ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] root node: %s(%p)\n",
+      ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Root node: %s(%p)\n",
               GetGraphNodeTypeString(entry->GetType()), entry);
     }
   }
@@ -388,7 +388,7 @@ void Graph::GetRunList(std::vector<std::vector<Node>>& parallelLists,
   }
   for (size_t i = 0; i < parallelLists.size(); i++) {
     for (size_t j = 0; j < parallelLists[i].size(); j++) {
-      ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] list %d - %s(%p)\n", i + 1,
+      ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] List %d - %s(%p)\n", i + 1,
               GetGraphNodeTypeString(parallelLists[i][j]->GetType()), parallelLists[i][j]);
     }
   }

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -35,6 +35,11 @@
 #include "hip_mempool_impl.hpp"
 #include "hip_vm.hpp"
 
+typedef struct ihipExtKernelEvents {
+  hipEvent_t startEvent_;
+  hipEvent_t stopEvent_;
+} ihipExtKernelEvents;
+
 namespace hip {
 struct Graph;
 struct GraphNode;
@@ -741,6 +746,7 @@ class GraphKernelNode : public GraphNode {
   unsigned int numParams_;
   hipKernelNodeAttrValue kernelAttr_;
   unsigned int kernelAttrInUse_;
+  ihipExtKernelEvents kernelEvents_;
 
  public:
     void PrintAttributes(std::ostream& out, hipGraphDebugDotFlags flag) {
@@ -823,8 +829,6 @@ class GraphKernelNode : public GraphNode {
       // capturehipExtModuleLaunchKernel() mixes host function with hipFunction_t, so we convert
       // here. If it's wrong, later functions will fail
       func = static_cast<hipFunction_t>(params.func);
-      ClPrint(amd::LOG_INFO, amd::LOG_CODE,
-              "[hipGraph] capturehipExtModuleLaunchKernel() should be called", status);
     } else if (status != hipSuccess) {
       ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] getStatFunc() failed with err %d", status);
     }
@@ -887,9 +891,13 @@ class GraphKernelNode : public GraphNode {
     return hipSuccess;
   }
 
-  GraphKernelNode(const hipKernelNodeParams* pNodeParams)
+  GraphKernelNode(const hipKernelNodeParams* pNodeParams, const ihipExtKernelEvents* pEvents)
       : GraphNode(hipGraphNodeTypeKernel, "bold", "octagon", "KERNEL") {
     kernelParams_ = *pNodeParams;
+    kernelEvents_ = { 0 };
+    if (pEvents != nullptr) {
+      kernelEvents_ = *pEvents;
+    }
     if (copyParams(pNodeParams) != hipSuccess) {
       ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to copy params");
     }
@@ -923,6 +931,7 @@ class GraphKernelNode : public GraphNode {
 
   GraphKernelNode(const GraphKernelNode& rhs) : GraphNode(rhs) {
     kernelParams_ = rhs.kernelParams_;
+    kernelEvents_ = rhs.kernelEvents_;
     hipError_t status = copyParams(&rhs.kernelParams_);
     if (status != hipSuccess) {
       ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to allocate memory to copy params");
@@ -957,8 +966,8 @@ class GraphKernelNode : public GraphNode {
         kernelParams_.gridDim.y * kernelParams_.blockDim.y,
         kernelParams_.gridDim.z * kernelParams_.blockDim.z, kernelParams_.blockDim.x,
         kernelParams_.blockDim.y, kernelParams_.blockDim.z, kernelParams_.sharedMemBytes,
-        stream, kernelParams_.kernelParams, kernelParams_.extra, nullptr, nullptr, 0, 0, 0, 0, 0,
-        0, 0);
+        stream, kernelParams_.kernelParams, kernelParams_.extra, kernelEvents_.startEvent_,
+        kernelEvents_.stopEvent_, 0, 0, 0, 0, 0, 0, 0);
     commands_.emplace_back(command);
     return status;
   }
@@ -1472,7 +1481,7 @@ class GraphMemcpyNodeFromSymbol : public GraphMemcpyNode1D {
     amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
     if (dstMemory == nullptr && kind != hipMemcpyDeviceToHost && kind != hipMemcpyDefault) {
       return hipErrorInvalidMemcpyDirection;
-    } else if (dstMemory != nullptr && dstMemory->getMemFlags() == 0 && 
+    } else if (dstMemory != nullptr && dstMemory->getMemFlags() == 0 &&
                kind != hipMemcpyDeviceToDevice && kind != hipMemcpyDefault) {
       return hipErrorInvalidMemcpyDirection;
     } else if (kind == hipMemcpyHostToHost || kind == hipMemcpyHostToDevice) {
@@ -1766,7 +1775,7 @@ class GraphEventRecordNode : public GraphNode {
       hipError_t status = e->enqueueRecordCommand(stream, commands_[0], true);
       if (status != hipSuccess) {
         ClPrint(amd::LOG_ERROR, amd::LOG_CODE,
-                "[hipGraph] enqueue event record command failed for node %p - status %d\n", this,
+                "[hipGraph] Enqueue event record command failed for node %p - status %d\n", this,
                 status);
       }
     }
@@ -1818,7 +1827,7 @@ class GraphEventWaitNode : public GraphNode {
       hipError_t status = e->enqueueStreamWaitCommand(stream, commands_[0]);
       if (status != hipSuccess) {
         ClPrint(amd::LOG_ERROR, amd::LOG_CODE,
-                "[hipGraph] enqueue stream wait command failed for node %p - status %d\n", this,
+                "[hipGraph] Enqueue stream wait command failed for node %p - status %d\n", this,
                 status);
       }
       commands_[0]->release();

projects/clr/hipamd/src/hip_stream.cpp

@@ -511,9 +511,6 @@ void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_d
 
 // ================================================================================================
 hipError_t hipStreamWaitEvent_common(hipStream_t stream, hipEvent_t event, unsigned int flags) {
-  ClPrint(amd::LOG_INFO, amd::LOG_API,
-          "[hipGraph] current capture node StreamWaitEvent on stream : %p, Event %p", stream,
-          event);
   hipError_t status = hipSuccess;
   if (event == nullptr || !hip::isValid(stream)) {
     return hipErrorInvalidHandle;
@@ -523,6 +520,9 @@ hipError_t hipStreamWaitEvent_common(hipStream_t stream, hipEvent_t event, unsig
   hip::Stream* eventStream = reinterpret_cast<hip::Stream*>(e->GetCaptureStream());
 
   if (eventStream != nullptr && eventStream->IsEventCaptured(event) == true) {
+    ClPrint(amd::LOG_INFO, amd::LOG_API,
+          "[hipGraph] Current capture node StreamWaitEvent on stream : %p, Event %p", stream,
+          event);
     if (waitStream == nullptr) {
       return hipErrorInvalidHandle;
     }