SWDEV-493526 - Create kernel node when hipLaunchByPtr is captured

Change-Id: Id3493485dfdb468436ab33e6d7cb19b6b0066fd4


[ROCm/clr commit: e08df57502]

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<hip::Stream*>(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);

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

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