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)); }