diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 5c746f1dc0..745f6f8101 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/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 lastCapturedNodes = s->GetLastCapturedNodes(); if (!lastCapturedNodes.empty()) { diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index c0567381e4..aea9a98482 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/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(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(const_cast(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(graph), - reinterpret_cast(pDependencies), numDependencies, pNodeParams, false); + reinterpret_cast(pDependencies), numDependencies, pNodeParams, + nullptr, false); *pGraphNode = reinterpret_cast(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(graph), - reinterpret_cast(pDependencies), numDependencies, pCopyParams, false); + reinterpret_cast(pDependencies), numDependencies, pCopyParams, + false); *pGraphNode = reinterpret_cast(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(graph), - reinterpret_cast(pDependencies), - numDependencies, dst, src, count, kind, false); + reinterpret_cast(pDependencies), + numDependencies, dst, src, count, kind, false); *pGraphNode = reinterpret_cast(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(clonedNode)->SetParams(dst, src, count, kind)); + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(dst, src, + count, kind)); } hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -1172,7 +1179,7 @@ hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, ihipGraphAddMemsetNode(&node, reinterpret_cast(graph), reinterpret_cast(pDependencies), numDependencies, pMemsetParams, false); - *pGraphNode = reinterpret_cast(node); + *pGraphNode = reinterpret_cast(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(node)) || pNodeParams == nullptr || pNodeParams->func == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || + pNodeParams == nullptr || pNodeParams->func == nullptr) { HIP_RETURN(hipErrorInvalidValue); } HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); @@ -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(node)) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || + pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } reinterpret_cast(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(node)) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || + pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } HIP_RETURN(reinterpret_cast(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(node); - if (hGraphExec == nullptr || !hip::GraphNode::isNodeValid(reinterpret_cast(n))) { + if (hGraphExec == nullptr || + !hip::GraphNode::isNodeValid(reinterpret_cast(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(node)) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || + pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } reinterpret_cast(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(node)) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(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> edges = reinterpret_cast(graph)->GetEdges(); + const std::vector> edges = + reinterpret_cast(graph)->GetEdges(); // returns only the number of edges in numEdges when from and to are null if (from == nullptr && to == nullptr) { *numEdges = edges.size(); @@ -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(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(graph)) || (pDependencies == nullptr && numDependencies > 0)) { + !hip::Graph::isGraphValid(reinterpret_cast(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(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(node)) || pNodeParams == nullptr) { + if (!hip::GraphNode::isNodeValid(reinterpret_cast(node)) || + pNodeParams == nullptr) { HIP_RETURN(hipErrorInvalidValue); } reinterpret_cast(node)->GetParams(pNodeParams); diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index cb71e4583f..9dba23ae1c 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -285,7 +285,7 @@ std::vector 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>& 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]); } } diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 3cf402c24b..9038ce937e 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/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(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(); diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index fa471f3ee5..8ee1ce9e9e 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/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(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; }