SWDEV-354717 - Add support for hipExtLaunchKernel as part of stream capture

Change-Id: I22708923f454a0f4456ff99d25559daffe08c208


[ROCm/clr commit: 2b322a9087]

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<hipFunction_t>(const_cast<void*>(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,

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,

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