From 0d6e8e378ee42702fcfd7e7e11da818031f64cb3 Mon Sep 17 00:00:00 2001 From: Anusha Godavarthy Surya Date: Tue, 2 Aug 2022 04:37:25 -0700 Subject: [PATCH] SWDEV-350474 - Added hipExtModuleLaunchKernel API support as part of stream capture Change-Id: I90a880ae0d3a85a0cc8380d2cb21e4759ea8151b --- hipamd/src/hip_graph.cpp | 69 ++++++++++++++++++++++++++++---- hipamd/src/hip_graph_capture.hpp | 8 ++++ hipamd/src/hip_module.cpp | 4 ++ 3 files changed, 74 insertions(+), 7 deletions(-) diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 007779ccb6..d081c584b5 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -70,13 +70,20 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const hipKernelNodeParams* pNodeParams) { if (pGraphNode == nullptr || graph == nullptr || (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr || - pNodeParams->func == nullptr || pNodeParams->kernelParams == nullptr) { + pNodeParams->func == nullptr) { return hipErrorInvalidValue; } + if (!ihipGraph::isGraphValid(graph)) { return hipErrorInvalidValue; } + // If neither 'kernelParams' or 'extra' are provided or if both are provided, return error + if ((pNodeParams->kernelParams == nullptr && pNodeParams->extra == nullptr) || + (pNodeParams->kernelParams != nullptr && pNodeParams->extra != nullptr)) { + return hipErrorInvalidValue; + } + hipFunction_t func = nullptr; hipError_t status = hipGraphKernelNode::getFunc(&func, *pNodeParams, ihipGetDevice()); if (status != hipSuccess) { @@ -88,12 +95,6 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, return status; } - // If neither 'kernelParams' or 'extra' are provided or if both are provided, return error - if ((pNodeParams->kernelParams == nullptr && pNodeParams->extra == nullptr) || - (pNodeParams->kernelParams != nullptr) && (pNodeParams->extra != nullptr)) { - return hipErrorInvalidValue; - } - size_t globalWorkSizeX = static_cast(pNodeParams->gridDim.x) * pNodeParams->blockDim.x; size_t globalWorkSizeY = static_cast(pNodeParams->gridDim.y) * pNodeParams->blockDim.y; size_t globalWorkSizeZ = static_cast(pNodeParams->gridDim.z) * pNodeParams->blockDim.z; @@ -204,6 +205,59 @@ hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction return hipSuccess; } +hipError_t capturehipExtModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, + uint32_t& globalWorkSizeX, uint32_t& globalWorkSizeY, + uint32_t& globalWorkSizeZ, uint32_t& localWorkSizeX, + uint32_t& localWorkSizeY, uint32_t& localWorkSizeZ, + size_t& sharedMemBytes, void**& kernelParams, + void**& extra, hipEvent_t& startEvent, + hipEvent_t& stopEvent, uint32_t& flags) { + ClPrint(amd::LOG_INFO, amd::LOG_API, + "[hipGraph] current capture node Ext kernel launch on stream : %p", stream); + if (!hip::isValid(stream)) { + return hipErrorInvalidValue; + } + hip::Stream* s = reinterpret_cast(stream); + + hipGraphNode_t pGraphNode; + hipError_t status; + if (startEvent != nullptr) { + pGraphNode = new hipGraphEventRecordNode(startEvent); + status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size()); + if (status != hipSuccess) { + return status; + } + s->SetLastCapturedNode(pGraphNode); + } + hipKernelNodeParams nodeParams; + nodeParams.func = f; + nodeParams.blockDim = dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ); + nodeParams.extra = extra; + nodeParams.gridDim = dim3(globalWorkSizeX / localWorkSizeX, globalWorkSizeY / localWorkSizeY, + globalWorkSizeZ / localWorkSizeZ); + nodeParams.kernelParams = kernelParams; + nodeParams.sharedMemBytes = sharedMemBytes; + status = + ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &nodeParams); + + if (status != hipSuccess) { + return status; + } + s->SetLastCapturedNode(pGraphNode); + if (stopEvent != nullptr) { + pGraphNode = new hipGraphEventRecordNode(stopEvent); + status = ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size()); + if (status != hipSuccess) { + return status; + } + s->SetLastCapturedNode(pGraphNode); + } + return hipSuccess; +} + hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*& p) { ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy3D on stream : %p", stream); @@ -910,6 +964,7 @@ hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, (numDependencies > 0 && pDependencies == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } + HIP_RETURN_DURATION( ihipGraphAddKernelNode(pGraphNode, graph, pDependencies, numDependencies, pNodeParams)); } diff --git a/hipamd/src/hip_graph_capture.hpp b/hipamd/src/hip_graph_capture.hpp index 0ba06f3a72..d09041b26f 100644 --- a/hipamd/src/hip_graph_capture.hpp +++ b/hipamd/src/hip_graph_capture.hpp @@ -23,6 +23,14 @@ hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction, dim3& gridDim, dim3& blockDim, void**& args, size_t& sharedMemBytes); +hipError_t capturehipExtModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, + uint32_t& globalWorkSizeX, uint32_t& globalWorkSizeY, + uint32_t& globalWorkSizeZ, uint32_t& localWorkSizeX, + uint32_t& localWorkSizeY, uint32_t& localWorkSizeZ, + size_t& sharedMemBytes, void**& kernelParams, + void**& extra, hipEvent_t& startEvent, + hipEvent_t& stopEvent, uint32_t& flags); + 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); diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 7921de3ac9..fde15a4a18 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -420,6 +420,10 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags); + STREAM_CAPTURE(hipExtModuleLaunchKernel, hStream, f, globalWorkSizeX, globalWorkSizeY, + globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, + kernelParams, extra, startEvent, stopEvent, flags); + HIP_RETURN(ihipModuleLaunchKernel(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags));