From ed2f9051fa0fdc3f602bbe1f933ce53c8b185046 Mon Sep 17 00:00:00 2001 From: Anusha Godavarthy Surya Date: Thu, 1 Sep 2022 05:01:55 +0000 Subject: [PATCH] SWDEV-354717 - Add support for hipExtLaunchKernel as part of stream capture Change-Id: I22708923f454a0f4456ff99d25559daffe08c208 [ROCm/clr commit: 2b322a90873955ed7aa02d0f7b7537abb8c2f045] --- projects/clr/hipamd/src/hip_graph.cpp | 40 ++++++++++++++----- projects/clr/hipamd/src/hip_graph_capture.hpp | 4 ++ projects/clr/hipamd/src/hip_module.cpp | 2 + 3 files changed, 37 insertions(+), 9 deletions(-) diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 0246742758..edf6ee97c8 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -205,15 +205,12 @@ 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); +hipError_t ihipExtLaunchKernel(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) { if (!hip::isValid(stream)) { return hipErrorInvalidValue; } @@ -258,6 +255,31 @@ hipError_t capturehipExtModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f 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 Module launch kernel on stream : %p", stream); + return ihipExtLaunchKernel(stream, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, + localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, + kernelParams, extra, startEvent, stopEvent, flags); +} + +hipError_t capturehipExtLaunchKernel(hipStream_t& stream, const void*& hostFunction, dim3& gridDim, + 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); + return ihipExtLaunchKernel(stream, + reinterpret_cast(const_cast(hostFunction)), + gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, + sharedMemBytes, args, nullptr, startEvent, stopEvent, flags); +} + hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, uint32_t& gridDimX, uint32_t& gridDimY, uint32_t& gridDimZ, uint32_t& blockDimX, uint32_t& blockDimY, uint32_t& blockDimZ, diff --git a/projects/clr/hipamd/src/hip_graph_capture.hpp b/projects/clr/hipamd/src/hip_graph_capture.hpp index 6650ec9a93..9fbcd01224 100644 --- a/projects/clr/hipamd/src/hip_graph_capture.hpp +++ b/projects/clr/hipamd/src/hip_graph_capture.hpp @@ -31,6 +31,10 @@ hipError_t capturehipExtModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f void**& extra, hipEvent_t& startEvent, hipEvent_t& stopEvent, uint32_t& flags); +hipError_t capturehipExtLaunchKernel(hipStream_t& stream, const void*& hostFunction, dim3& gridDim, + dim3& blockDim, void**& args, size_t& sharedMemBytes, + hipEvent_t& startEvent, hipEvent_t& stopEvent, int& flags); + hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, uint32_t& gridDimX, uint32_t& gridDimY, uint32_t& gridDimZ, uint32_t& blockDimX, uint32_t& blockDimY, uint32_t& blockDimZ, diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 69abc42193..ade9529844 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -487,6 +487,8 @@ extern "C" hipError_t hipExtLaunchKernel(const void* hostFunction, dim3 gridDim, void** args, size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent, hipEvent_t stopEvent, int flags) { HIP_INIT_API(hipExtLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); + STREAM_CAPTURE(hipExtLaunchKernel, stream, hostFunction, gridDim, blockDim, args, sharedMemBytes, + startEvent, stopEvent, flags); HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, startEvent, stopEvent, flags)); }