SWDEV-240807 - Added Implementation for few more graph APIs

Change-Id: I76336a22233a208a3f54ff9e90f0c5bf4a1bddb4

hipamd/include/hip/amd_detail/hip_prof_str.h

@@ -236,14 +236,23 @@ 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_hipGraphAddDependencies = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphAddKernelNode = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphAddMemcpyNode = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphAddMemsetNode = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphCreate = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphDestroy = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphExecDestroy = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphGetNodes = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphGetRootNodes = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphInstantiate = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphKernelNodeGetParams = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphKernelNodeSetParams = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphLaunch = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphMemcpyNodeGetParams = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphMemcpyNodeSetParams = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphMemsetNodeGetParams = HIP_API_ID_NONE,
+  HIP_API_ID_hipGraphMemsetNodeSetParams = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphicsGLRegisterBuffer = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphicsMapResources = HIP_API_ID_NONE,
   HIP_API_ID_hipGraphicsResourceGetMappedPointer = HIP_API_ID_NONE,
@@ -3136,6 +3145,8 @@ typedef struct hip_api_data_s {
 #define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {};
 // hipGetTextureReference()
 #define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {};
+// hipGraphAddDependencies()
+#define INIT_hipGraphAddDependencies_CB_ARGS_DATA(cb_data) {};
 // hipGraphAddKernelNode()
 #define INIT_hipGraphAddKernelNode_CB_ARGS_DATA(cb_data) {};
 // hipGraphAddMemcpyNode()
@@ -3148,10 +3159,26 @@ typedef struct hip_api_data_s {
 #define INIT_hipGraphDestroy_CB_ARGS_DATA(cb_data) {};
 // hipGraphExecDestroy()
 #define INIT_hipGraphExecDestroy_CB_ARGS_DATA(cb_data) {};
+// hipGraphGetNodes()
+#define INIT_hipGraphGetNodes_CB_ARGS_DATA(cb_data) {};
+// hipGraphGetRootNodes()
+#define INIT_hipGraphGetRootNodes_CB_ARGS_DATA(cb_data) {};
 // hipGraphInstantiate()
 #define INIT_hipGraphInstantiate_CB_ARGS_DATA(cb_data) {};
+// hipGraphKernelNodeGetParams()
+#define INIT_hipGraphKernelNodeGetParams_CB_ARGS_DATA(cb_data) {};
+// hipGraphKernelNodeSetParams()
+#define INIT_hipGraphKernelNodeSetParams_CB_ARGS_DATA(cb_data) {};
 // hipGraphLaunch()
 #define INIT_hipGraphLaunch_CB_ARGS_DATA(cb_data) {};
+// hipGraphMemcpyNodeGetParams()
+#define INIT_hipGraphMemcpyNodeGetParams_CB_ARGS_DATA(cb_data) {};
+// hipGraphMemcpyNodeSetParams()
+#define INIT_hipGraphMemcpyNodeSetParams_CB_ARGS_DATA(cb_data) {};
+// hipGraphMemsetNodeGetParams()
+#define INIT_hipGraphMemsetNodeGetParams_CB_ARGS_DATA(cb_data) {};
+// hipGraphMemsetNodeSetParams()
+#define INIT_hipGraphMemsetNodeSetParams_CB_ARGS_DATA(cb_data) {};
 // hipGraphicsGLRegisterBuffer()
 #define INIT_hipGraphicsGLRegisterBuffer_CB_ARGS_DATA(cb_data) {};
 // hipGraphicsMapResources()

hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h

@@ -2292,6 +2292,64 @@ inline static hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGr
         cudaGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, pMemsetParams));
 }
 
+inline static hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes,
+                                          size_t* numNodes) {
+    return hipCUDAErrorTohipError(cudaGraphGetNodes(graph, nodes, numNodes));
+}
+
+inline static hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes,
+                                              size_t* pNumRootNodes) {
+    return hipCUDAErrorTohipError(cudaGraphGetRootNodes(graph, pRootNodes, pNumRootNodes));
+}
+
+inline static hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node,
+                                                     hipKernelNodeParams* pNodeParams) {
+    return hipCUDAErrorTohipError(cudaGraphKernelNodeGetParams(node, pNodeParams));
+}
+
+inline static hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node,
+                                                     const hipKernelNodeParams* pNodeParams) {
+    return hipCUDAErrorTohipError(cudaGraphKernelNodeSetParams(node, pNodeParams));
+}
+
+inline static hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node,
+                                                     hipMemcpy3DParms* pNodeParams) {
+    return hipCUDAErrorTohipError(cudaGraphMemcpyNodeGetParams(node, pNodeParams));
+}
+
+inline static hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node,
+                                                     const hipMemcpy3DParms* pNodeParams) {
+    return hipCUDAErrorTohipError(cudaGraphMemcpyNodeSetParams(node, pNodeParams));
+}
+
+inline static hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node,
+                                                     hipMemsetParams* pNodeParams) {
+    return hipCUDAErrorTohipError(cudaGraphMemsetNodeGetParams(node, pNodeParams));
+}
+
+inline static hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node,
+                                                     const hipMemsetParams* pNodeParams) {
+    return hipCUDAErrorTohipError(cudaGraphMemsetNodeSetParams(node, pNodeParams));
+}
+
+inline static hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec,
+                                                         hipGraphNode_t node,
+                                                         const hipKernelNodeParams* pNodeParams) {
+    return hipCUDAErrorTohipError(cudaGraphExecKernelNodeSetParams(hGraphExec, node, pNodeParams));
+}
+
+inline static hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
+                                                 const hipGraphNode_t* to, size_t numDependencies) {
+    return hipCUDAErrorTohipError(cudaGraphAddDependencies(graph, from, to, numDependencies));
+}
+
+inline static hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                              const hipGraphNode_t* pDependencies,
+                                              size_t numDependencies) {
+    return hipCUDAErrorTohipError(
+      cudaGraphAddEmptyNode(pGraphNode, graph, pDependencies, numDependencies));
+}
+
 inline static hipError_t hipStreamWriteValue32(hipStream_t stream,
                                                void* ptr, int32_t value, unsigned int flags) {
    if (value < 0) {

hipamd/src/amdhip.def

@@ -291,3 +291,14 @@ hipImportExternalSemaphore
 hipSignalExternalSemaphoresAsync
 hipWaitExternalSemaphoresAsync
 hipDestroyExternalSemaphore
+hipGraphGetNodes
+hipGraphGetRootNodes
+hipGraphKernelNodeGetParams
+hipGraphKernelNodeSetParams
+hipGraphMemcpyNodeGetParams
+hipGraphMemcpyNodeSetParams
+hipGraphMemsetNodeGetParams
+hipGraphMemsetNodeSetParams
+hipGraphExecKernelNodeSetParams
+hipGraphAddDependencies
+hipGraphAddEmptyNode

hipamd/src/hip_graph.cpp

@@ -413,3 +413,100 @@ hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) {
   HIP_INIT_API(hipGraphLaunch, graphExec, stream);
   HIP_RETURN_DURATION(ihipGraphlaunch(graphExec, 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();
+  }
+  if (*numNodes > 0) {
+    nodes = graph->GetNodes().data();
+  }
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes,
+                                size_t* pNumRootNodes) {
+  HIP_INIT_API(hipGraphGetRootNodes, graph, pRootNodes, pNumRootNodes);
+  if (graph == nullptr || pNumRootNodes == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  std::vector<Node> rootNodes = graph->GetRootNodes();
+  pRootNodes = rootNodes.data();
+  *pNumRootNodes = rootNodes.size();
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams) {
+  HIP_INIT_API(hipGraphKernelNodeGetParams, node, pNodeParams);
+  if (node == nullptr || pNodeParams == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<hipGraphKernelNode*>(node)->GetParams(pNodeParams);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node,
+                                       const hipKernelNodeParams* pNodeParams) {
+  HIP_INIT_API(hipGraphKernelNodeSetParams, node, pNodeParams);
+  if (node == nullptr || pNodeParams == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<hipGraphKernelNode*>(node)->SetParams(pNodeParams);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams) {
+  HIP_INIT_API(hipGraphMemcpyNodeGetParams, node, pNodeParams);
+  if (node == nullptr || pNodeParams == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<hipGraphMemcpyNode*>(node)->GetParams(pNodeParams);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams) {
+  HIP_INIT_API(hipGraphMemcpyNodeSetParams, node, pNodeParams);
+  if (node == nullptr || pNodeParams == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<hipGraphMemcpyNode*>(node)->SetParams(pNodeParams);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams) {
+  HIP_INIT_API(hipGraphMemsetNodeGetParams, node, pNodeParams);
+  if (node == nullptr || pNodeParams == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<hipGraphMemsetNode*>(node)->GetParams(pNodeParams);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams) {
+  HIP_INIT_API(hipGraphMemsetNodeSetParams, node, pNodeParams);
+  if (node == nullptr || pNodeParams == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  reinterpret_cast<hipGraphMemsetNode*>(node)->SetParams(pNodeParams);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
+                                   const hipGraphNode_t* to, size_t numDependencies) {
+  HIP_INIT_API(hipGraphAddDependencies, graph, from, to, numDependencies);
+  if (graph == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  if (numDependencies == 0) {
+    HIP_RETURN(hipSuccess);
+  } else if (from == nullptr || to == nullptr) {
+    return HIP_RETURN(hipErrorInvalidValue);
+  }
+  for (size_t i = 0; i < numDependencies; i++) {
+    if (graph->AddEdge(from[i], to[i]) != hipSuccess) {
+      HIP_RETURN(hipErrorInvalidValue);
+    }
+  }
+  HIP_RETURN(hipSuccess);
+}

hipamd/src/hip_graph_internal.hpp

@@ -120,7 +120,7 @@ class hipGraphKernelNode : public hipGraphNode {
   void GetParams(hipKernelNodeParams* params) {
     std::memcpy(params, pKernelParams_, sizeof(hipKernelNodeParams));
   }
-  void SetParams(hipKernelNodeParams* params) {
+  void SetParams(const hipKernelNodeParams* params) {
     std::memcpy(pKernelParams_, params, sizeof(hipKernelNodeParams));
   }
 };
@@ -145,7 +145,7 @@ class hipGraphMemcpyNode : public hipGraphNode {
   void GetParams(hipMemcpy3DParms* params) {
     std::memcpy(params, pCopyParams_, sizeof(hipMemcpy3DParms));
   }
-  void SetParams(hipMemcpy3DParms* params) {
+  void SetParams(const hipMemcpy3DParms* params) {
     std::memcpy(pCopyParams_, params, sizeof(hipMemcpy3DParms));
   }
 };
@@ -263,7 +263,7 @@ class hipGraphMemsetNode : public hipGraphNode {
   void GetParams(hipMemsetParams* params) {
     std::memcpy(params, pMemsetParams_, sizeof(hipMemsetParams));
   }
-  void SetParams(hipMemsetParams* params) {
+  void SetParams(const hipMemsetParams* params) {
     std::memcpy(pMemsetParams_, params, sizeof(hipMemsetParams));
   }
 };

hipamd/src/hip_hcc.def.in

@@ -297,3 +297,14 @@ hipGraphicsMapResources
 hipGraphicsResourceGetMappedPointer
 hipGraphicsUnmapResources
 hipGraphicsUnregisterResource
+hipGraphGetNodes
+hipGraphGetRootNodes
+hipGraphKernelNodeGetParams
+hipGraphKernelNodeSetParams
+hipGraphMemcpyNodeGetParams
+hipGraphMemcpyNodeSetParams
+hipGraphMemsetNodeGetParams
+hipGraphMemsetNodeSetParams
+hipGraphExecKernelNodeSetParams
+hipGraphAddDependencies
+hipGraphAddEmptyNode

hipamd/src/hip_hcc.map.in

@@ -314,3 +314,20 @@ global:
 local:
     *;
 } hip_4.2;
+
+hip_4.4 {
+global:
+    hipGraphGetNodes;
+    hipGraphGetRootNodes;
+    hipGraphKernelNodeGetParams;
+    hipGraphKernelNodeSetParams;
+    hipGraphMemcpyNodeGetParams;
+    hipGraphMemcpyNodeSetParams;
+    hipGraphMemsetNodeGetParams;
+    hipGraphMemsetNodeSetParams;
+    hipGraphExecKernelNodeSetParams;
+    hipGraphAddDependencies;
+    hipGraphAddEmptyNode;
+local:
+    *;
+}hip_4.3;