diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index e63730aa26..49ea1fbe94 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -348,6 +348,35 @@ hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, u return hipSuccess; } +hipError_t capturehipLaunchByPtr(hipStream_t& stream, hipFunction_t func, dim3 blockDim, + dim3 gridDim, unsigned int sharedMemBytes, void** extra) { + ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node LaunchByPtr on stream : %p", + stream); + if (!hip::isValid(stream)) { + return hipErrorContextIsDestroyed; + } + + hipKernelNodeParams nodeParams; + nodeParams.func = func; + nodeParams.blockDim = blockDim; + nodeParams.gridDim = gridDim; + nodeParams.sharedMemBytes = sharedMemBytes; + nodeParams.extra = extra; + nodeParams.kernelParams = nullptr; + + hip::GraphNode* pGraphNode; + hip::Stream* s = reinterpret_cast(stream); + hipError_t status = + ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &nodeParams); + 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); diff --git a/projects/clr/hipamd/src/hip_graph_capture.hpp b/projects/clr/hipamd/src/hip_graph_capture.hpp index f1d96898e4..5c02648f99 100644 --- a/projects/clr/hipamd/src/hip_graph_capture.hpp +++ b/projects/clr/hipamd/src/hip_graph_capture.hpp @@ -42,6 +42,9 @@ hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, u uint32_t& sharedMemBytes, void**& kernelParams, void**& extra); +hipError_t capturehipLaunchByPtr(hipStream_t& stream, hipFunction_t func, dim3 blockDim, + dim3 gridDim, unsigned int sharedMemBytes, void** extra); + 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/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp index 5f0d355172..9c0c99d6ab 100644 --- a/projects/clr/hipamd/src/hip_platform.cpp +++ b/projects/clr/hipamd/src/hip_platform.cpp @@ -279,6 +279,9 @@ hipError_t hipLaunchByPtr(const void* hostFunction) { void* extra[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec.arguments_[0], HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; + STREAM_CAPTURE(hipLaunchByPtr, exec.hStream_, func, exec.blockDim_, exec.gridDim_, + exec.sharedMem_, extra); + HIP_RETURN(hipModuleLaunchKernel(func, exec.gridDim_.x, exec.gridDim_.y, exec.gridDim_.z, exec.blockDim_.x, exec.blockDim_.y, exec.blockDim_.z, exec.sharedMem_, exec.hStream_, nullptr, extra));