SWDEV-240806 - Added API hipGraphAddMemcpyNode1D, hipGraphAddEmptyNode, hipGraphExecKernelNodeSetParams

Change-Id: I0d7ec8c0ea1abc3fc0f1e10fa7865f355d9cf2ad


[ROCm/clr commit: a71f54d9e9]

projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h

@@ -279,6 +279,7 @@ enum hip_api_id_t {
   HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NONE,
   HIP_API_ID_hipGetTextureObjectTextureDesc = HIP_API_ID_NONE,
   HIP_API_ID_hipGetTextureReference = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphAddMemcpyNode1D = HIP_API_ID_NONE,
   HIP_API_ID_hipMemcpy2DArrayToArray = HIP_API_ID_NONE,
   HIP_API_ID_hipMemcpyArrayToArray = HIP_API_ID_NONE,
   HIP_API_ID_hipMemcpyAtoA = HIP_API_ID_NONE,
@@ -3769,6 +3770,8 @@ typedef struct hip_api_data_s {
 #define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {};
 // hipGetTextureReference()
 #define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {};
+// hipGraphAddMemcpyNode1D()
+#define INIT_hipGraphAddMemcpyNode1D_CB_ARGS_DATA(cb_data) {};
 // hipMemcpy2DArrayToArray()
 #define INIT_hipMemcpy2DArrayToArray_CB_ARGS_DATA(cb_data) {};
 // hipMemcpyArrayToArray()

projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h

@@ -2292,6 +2292,13 @@ inline static hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGr
         cudaGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, pCopyParams));
 }
 
+inline static hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                   const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                   void* dst, const void* src, size_t count, hipMemcpyKind kind) {
+    return hipCUDAErrorTohipError(
+        cudaGraphAddMemcpyNode1D(pGraphNode, graph, pDependencies, numDependencies, dst, src, count, kind));
+}
+
 inline static hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                                const hipGraphNode_t* pDependencies,
                                                size_t numDependencies,

projects/clr/hipamd/src/amdhip.def

@@ -281,6 +281,7 @@ hipGraphDestroy
 hipGraphAddKernelNode
 hipGraphAddMemsetNode
 hipGraphAddMemcpyNode
+hipGraphAddMemcpyNode1D
 hipGraphInstantiate
 hipGraphLaunch
 hipStreamIsCapturing
@@ -300,3 +301,5 @@ hipGraphMemcpyNodeSetParams
 hipGraphMemsetNodeGetParams
 hipGraphMemsetNodeSetParams
 hipGraphAddDependencies
+hipGraphExecKernelNodeSetParams
+hipGraphAddEmptyNode

projects/clr/hipamd/src/hip_graph.cpp

@@ -27,13 +27,7 @@
 thread_local std::vector<hipStream_t> g_captureStreams;
 std::unordered_map<amd::Command*, hipGraphExec_t> hipGraphExec::activeGraphExec_;
 
-hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
-                                  const hipGraphNode_t* pDependencies, size_t numDependencies,
-                                  const hipKernelNodeParams* pNodeParams) {
-  if (pGraphNode == nullptr || graph == nullptr ||
-      (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr) {
-    return hipErrorInvalidValue;
-  }
+hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams) {
   hipFunction_t func = nullptr;
   hipError_t status =
       PlatformState::instance().getStatFunc(&func, pNodeParams->func, ihipGetDevice());
@@ -56,6 +50,24 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
   if (status != hipSuccess) {
     return status;
   }
+}
+
+hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                  const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                  const hipKernelNodeParams* pNodeParams) {
+  if (pGraphNode == nullptr || graph == nullptr ||
+      (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr) {
+    return hipErrorInvalidValue;
+  }
+  hipError_t status = ihipValidateKernelParams(pNodeParams);
+  if (hipSuccess != status) {
+    return status;
+  }
+  hipFunction_t func = nullptr;
+  status = PlatformState::instance().getStatFunc(&func, pNodeParams->func, ihipGetDevice());
+  if ((status != hipSuccess) || (func == nullptr)) {
+    return hipErrorInvalidDeviceFunction;
+  }
   *pGraphNode = new hipGraphKernelNode(pNodeParams, func);
   if (numDependencies == 0) {
     graph->AddNode(*pGraphNode);
@@ -88,6 +100,26 @@ hipError_t ihipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
   return hipSuccess;
 }
 
+hipError_t ihipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                    const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                    void* dst, const void* src, size_t count, hipMemcpyKind kind) {
+  if (pGraphNode == nullptr || graph == nullptr ||
+      (numDependencies > 0 && pDependencies == nullptr)) {
+    return hipErrorInvalidValue;
+  }
+  ihipMemcpy_validate(dst, src, count, kind);
+  *pGraphNode = new hipGraphMemcpyNode1D(dst, src, count, kind);
+  if (numDependencies == 0) {
+    graph->AddNode(*pGraphNode);
+  }
+  for (size_t i = 0; i < numDependencies; i++) {
+    if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) {
+      return hipErrorInvalidValue;
+    }
+  }
+  return hipSuccess;
+}
+
 hipError_t ihipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                   const hipGraphNode_t* pDependencies, size_t numDependencies,
                                   const hipMemsetParams* pMemsetParams) {
@@ -268,6 +300,7 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns
   ClPrint(amd::LOG_INFO, amd::LOG_API,
           "[hipGraph] current capture node StreamWaitEvent on stream : %p, Event %p", stream,
           event);
+
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   hip::Event* e = reinterpret_cast<hip::Event*>(event);
 
@@ -281,7 +314,7 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns
   }
   s->AddCrossCapturedNode(e->GetNodesPrevToRecorded());
   g_captureStreams.push_back(stream);
-  return hipSuccess;
+  HIP_RETURN(hipSuccess);
 }
 
 hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus** pCaptureStatus) {
@@ -355,7 +388,7 @@ hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
   HIP_INIT_API(hipGraphAddKernelNode, pGraphNode, graph, pDependencies, numDependencies,
                pNodeParams);
   HIP_RETURN_DURATION(
-      ihipGraphAddKernelNode(pGraphNode, graph, pDependencies, numDependencies, pNodeParams););
+      ihipGraphAddKernelNode(pGraphNode, graph, pDependencies, numDependencies, pNodeParams));
 }
 
 hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
@@ -365,7 +398,17 @@ hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                pCopyParams);
 
   HIP_RETURN_DURATION(
-      ihipGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, pCopyParams););
+      ihipGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, pCopyParams));
+}
+
+hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                   const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                   void* dst, const void* src, size_t count, hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphAddMemcpyNode1D, pGraphNode, graph, pDependencies, numDependencies, dst, src,
+               count, kind);
+
+  HIP_RETURN_DURATION(ihipGraphAddMemcpyNode1D(pGraphNode, graph, pDependencies, numDependencies,
+                                               dst, src, count, kind));
 }
 
 hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
@@ -375,7 +418,26 @@ hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                pMemsetParams);
 
   HIP_RETURN_DURATION(
-      ihipGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, pMemsetParams););
+      ihipGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, pMemsetParams));
+}
+
+hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                const hipGraphNode_t* pDependencies, size_t numDependencies) {
+  HIP_INIT_API(hipGraphAddEmptyNode, pGraphNode, graph, pDependencies, numDependencies);
+  if (pGraphNode == nullptr || graph == nullptr ||
+      (numDependencies > 0 && pDependencies == nullptr)) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  *pGraphNode = new hipGraphEmptyNode();
+  if (numDependencies == 0) {
+    graph->AddNode(*pGraphNode);
+  }
+  for (size_t i = 0; i < numDependencies; i++) {
+    if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) {
+      return hipErrorInvalidValue;
+    }
+  }
+  return hipSuccess;
 }
 
 hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
@@ -417,8 +479,9 @@ hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) {
 hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes) {
   HIP_INIT_API(hipGraphGetNodes, graph, nodes, numNodes);
   if (graph == nullptr || numNodes == nullptr) {
-    *numNodes = graph->GetNodeCount();
+    return HIP_RETURN(hipErrorInvalidValue);
   }
+  *numNodes = graph->GetNodeCount();
   if (*numNodes > 0) {
     nodes = graph->GetNodes().data();
   }
@@ -510,3 +573,12 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
   }
   HIP_RETURN(hipSuccess);
 }
+
+hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
+                                           const hipKernelNodeParams* pNodeParams) {
+  HIP_INIT_API(hipGraphExecKernelNodeSetParams, hGraphExec, node, pNodeParams);
+  if (hGraphExec == nullptr || node == nullptr || pNodeParams == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  return reinterpret_cast<hipGraphKernelNode*>(node)->SetCommandParams(pNodeParams);
+}

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -32,6 +32,7 @@
 #include "hip_graph_helper.hpp"
 
 typedef hipGraphNode* Node;
+hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams);
 
 struct hipGraphNode {
  protected:
@@ -123,6 +124,24 @@ class hipGraphKernelNode : public hipGraphNode {
   void SetParams(const hipKernelNodeParams* params) {
     std::memcpy(pKernelParams_, params, sizeof(hipKernelNodeParams));
   }
+  hipError_t SetCommandParams(const hipKernelNodeParams* params) {
+    if (params->func != pKernelParams_->func) {
+      return hipErrorInvalidValue;
+    }
+    // updates kernel params
+    hipError_t status = ihipValidateKernelParams(params);
+    if (hipSuccess != status) {
+      return status;
+    }
+    size_t globalWorkOffset[3] = {0};
+    size_t globalWorkSize[3] = {params->gridDim.x, params->gridDim.y, params->gridDim.z};
+    size_t localWorkSize[3] = {params->blockDim.x, params->blockDim.y, params->blockDim.z};
+    reinterpret_cast<amd::NDRangeKernelCommand*>(commands_[0])
+        ->setSizes(globalWorkOffset, globalWorkSize, localWorkSize);
+    reinterpret_cast<amd::NDRangeKernelCommand*>(commands_[0])
+        ->setSharedMemBytes(params->sharedMemBytes);
+    return hipSuccess;
+  }
 };
 
 class hipGraphMemcpyNode : public hipGraphNode {
@@ -150,7 +169,6 @@ class hipGraphMemcpyNode : public hipGraphNode {
   }
 };
 
-
 class hipGraphMemcpyNode1D : public hipGraphNode {
   void* dst_;
   const void* src_;
@@ -315,6 +333,20 @@ class hipGraphHostNode : public hipGraphNode {
   }
 };
 
+class hipGraphEmptyNode : public hipGraphNode {
+ public:
+  hipGraphEmptyNode() : hipGraphNode(hipGraphNodeTypeEmpty) {}
+  ~hipGraphEmptyNode() {}
+
+  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;
+  }
+};
+
 struct hipGraphExec {
   std::vector<std::vector<Node>> parallelLists_;
   std::vector<Node> levelOrder_;

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

@@ -281,6 +281,7 @@ hipGraphDestroy
 hipGraphAddKernelNode
 hipGraphAddMemsetNode
 hipGraphAddMemcpyNode
+hipGraphAddMemcpyNode1D
 hipGraphInstantiate
 hipGraphLaunch
 hipStreamIsCapturing
@@ -306,3 +307,5 @@ hipGraphMemcpyNodeSetParams
 hipGraphMemsetNodeGetParams
 hipGraphMemsetNodeSetParams
 hipGraphAddDependencies
+hipGraphExecKernelNodeSetParams
+hipGraphAddEmptyNode

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

@@ -291,6 +291,7 @@ global:
     hipGraphAddKernelNode;
     hipGraphAddMemsetNode;
     hipGraphAddMemcpyNode;
+    hipGraphAddMemcpyNode1D;
     hipGraphInstantiate;
     hipGraphLaunch;
     hipStreamIsCapturing;
@@ -326,6 +327,8 @@ global:
     hipStreamWaitValue64;
     hipStreamWriteValue32;
     hipStreamWriteValue64;
+    hipGraphExecKernelNodeSetParams;
+    hipGraphAddEmptyNode;
 local:
     *;
 } hip_4.3;

projects/clr/hipamd/src/hip_stream.cpp

@@ -40,7 +40,8 @@ Stream::Stream(hip::Device* dev, Priority p, unsigned int f, bool null_stream,
       flags_(f),
       null_(null_stream),
       cuMask_(cuMask),
-      captureStatus_(captureStatus) {}
+      captureStatus_(captureStatus),
+      originStream_(false) {}
 
 // ================================================================================================
 Stream::~Stream() {