SWDEV-240806 - Added support for host graph node

Change-Id: I8b3088d34d0b4e8080b55b5da80d821c5f29d8a4


[ROCm/clr commit: ce58dae704]

projects/clr/hipamd/src/amdhip.def

@@ -340,3 +340,7 @@ hipGraphAddEventWaitNode
 hipGraphEventWaitNodeGetEvent
 hipGraphEventWaitNodeSetEvent
 hipGraphExecEventWaitNodeSetEvent
+hipGraphAddHostNode
+hipGraphHostNodeGetParams
+hipGraphHostNodeSetParams
+hipGraphExecHostNodeSetParams

projects/clr/hipamd/src/hip_graph.cpp

@@ -617,6 +617,23 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns
   return hipSuccess;
 }
 
+hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*& userData) {
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset2D on stream : %p",
+          stream);
+  if (fn == nullptr || userData == nullptr) {
+    return hipErrorInvalidValue;
+  }
+  hipHostNodeParams hostParams = {0};
+  hostParams.fn = fn;
+  hostParams.userData = userData;
+  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+  hipGraphNode_t pGraphNode = new hipGraphHostNode(&hostParams);
+  ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(),
+                   s->GetLastCapturedNodes().size());
+  s->SetLastCapturedNode(pGraphNode);
+  return hipSuccess;
+}
+
 hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) {
   HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus);
   if (stream == nullptr) {
@@ -1345,3 +1362,49 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph
   }
   HIP_RETURN(reinterpret_cast<hipGraphEventRecordNode*>(clonedNode)->SetParams(event));
 }
+
+hipError_t hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                               const hipGraphNode_t* pDependencies, size_t numDependencies,
+                               const hipHostNodeParams* pNodeParams) {
+  HIP_INIT_API(hipGraphAddHostNode, pGraphNode, graph, pDependencies, numDependencies, pNodeParams);
+  if (pGraphNode == nullptr || graph == nullptr ||
+      (numDependencies > 0 && pDependencies == nullptr)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  if (pNodeParams->fn == nullptr || pNodeParams->userData == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  *pGraphNode = new hipGraphHostNode(pNodeParams);
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNodeParams) {
+  HIP_INIT_API(hipGraphHostNodeGetParams, node, pNodeParams);
+  if (node == nullptr || pNodeParams == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<hipGraphHostNode*>(node)->GetParams(pNodeParams);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParams* pNodeParams) {
+  HIP_INIT_API(hipGraphHostNodeSetParams, node, pNodeParams);
+  if (pNodeParams->fn == nullptr || pNodeParams->userData == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  HIP_RETURN(reinterpret_cast<hipGraphHostNode*>(node)->SetParams(pNodeParams));
+}
+
+hipError_t hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
+                                         const hipHostNodeParams* pNodeParams) {
+  HIP_INIT_API(hipGraphExecHostNodeSetParams, hGraphExec, node, pNodeParams);
+  if (pNodeParams->fn == nullptr || pNodeParams->userData == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node);
+  if (clonedNode == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  HIP_RETURN(reinterpret_cast<hipGraphHostNode*>(clonedNode)->SetParams(pNodeParams));
+}

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -842,12 +842,42 @@ class hipGraphHostNode : public hipGraphNode {
     return new hipGraphHostNode(static_cast<hipGraphHostNode const&>(*this));
   }
 
-  hipError_t CreateCommand(amd::HostQueue* queue);
+  hipError_t CreateCommand(amd::HostQueue* queue) {
+    amd::Command::EventWaitList waitList;
+    commands_.reserve(1);
+    amd::Command* command = new amd::Marker(*queue, !kMarkerDisableFlush, waitList);
+    commands_.emplace_back(command);
+    return hipSuccess;
+  }
+
+  static void Callback(cl_event event, cl_int command_exec_status, void* user_data) {
+    hipHostNodeParams* pNodeParams = reinterpret_cast<hipHostNodeParams*>(user_data);
+    pNodeParams->fn(pNodeParams->userData);
+  }
+
+  void EnqueueCommands(hipStream_t stream) {
+    if (!commands_.empty()) {
+      if (!commands_[0]->setCallback(CL_COMPLETE, hipGraphHostNode::Callback, pNodeParams_)) {
+        ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed during setCallback");
+      }
+      commands_[0]->enqueue();
+      // Add the new barrier to stall the stream, until the callback is done
+      amd::Command::EventWaitList eventWaitList;
+      eventWaitList.push_back(commands_[0]);
+      amd::Command* block_command =
+          new amd::Marker(*commands_[0]->queue(), !kMarkerDisableFlush, eventWaitList);
+      if (block_command == nullptr) {
+        ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed during block command creation");
+      }
+      block_command->enqueue();
+      block_command->release();
+    }
+  }
 
   void GetParams(hipHostNodeParams* params) {
     std::memcpy(params, pNodeParams_, sizeof(hipHostNodeParams));
   }
-  hipError_t SetParams(hipHostNodeParams* params) {
+  hipError_t SetParams(const hipHostNodeParams* params) {
     std::memcpy(pNodeParams_, params, sizeof(hipHostNodeParams));
     return hipSuccess;
   }

projects/clr/hipamd/src/hip_hcc.def.in

@@ -338,3 +338,7 @@ hipGraphAddEventWaitNode
 hipGraphEventWaitNodeGetEvent
 hipGraphEventWaitNodeSetEvent
 hipGraphExecEventWaitNodeSetEvent
+hipGraphAddHostNode
+hipGraphHostNodeGetParams
+hipGraphHostNodeSetParams
+hipGraphExecHostNodeSetParams

projects/clr/hipamd/src/hip_hcc.map.in

@@ -370,6 +370,10 @@ global:
     hipGraphEventWaitNodeGetEvent;
     hipGraphEventWaitNodeSetEvent;
     hipGraphExecEventWaitNodeSetEvent;
+    hipGraphAddHostNode;
+    hipGraphHostNodeGetParams;
+    hipGraphHostNodeSetParams;
+    hipGraphExecHostNodeSetParams;
 local:
     *;
 } hip_4.4;