SWDEV-240806 - Initial commit for hipGraph and stream capture infrastructure

On StreamBegincapture captures the parameters passed to APIs and respective node will be created and added to graph
All parameters are passed to STREAM_CAPTURE macro, it checks if stream in capture mode and redirects the call to the capture function and returns
Updated hipStream and hipEvent with capture parameters
Added handling for hipStreamBeginCapture & hipStreamEndCapture

Change-Id: Ic8926a7b4336c2cc81f0b3a9a224aa392c474134


[ROCm/hip commit: c35ba37287]

projects/hip/include/hip/amd_detail/hip_runtime_api.h

@@ -3886,6 +3886,130 @@ const char* hipKernelNameRef(const hipFunction_t f);
 const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream);
 int hipGetStreamDeviceId(hipStream_t stream);
 
+#ifdef __cplusplus
+/**
+ * An opaque value that represents a hip graph
+ */
+class hipGraph;
+typedef hipGraph* hipGraph_t;
+
+/**
+ * An opaque value that represents a hip graph node
+ */
+class hipGraphNode;
+typedef hipGraphNode* hipGraphNode_t;
+
+/**
+ * An opaque value that represents a hip graph Exec
+ */
+class hipGraphExec;
+typedef hipGraphExec* hipGraphExec_t;
+typedef enum hipGraphNodeType {
+  hipGraphNodeTypeKernel = 1,             ///< GPU kernel node
+  hipGraphNodeTypeMemcpy = 2,             ///< Memcpy 3D node
+  hipGraphNodeTypeMemset = 3,             ///< Memset 1D node
+  hipGraphNodeTypeHost = 4,               ///< Host (executable) node
+  hipGraphNodeTypeGraph = 5,              ///< Node which executes an embedded graph
+  hipGraphNodeTypeEmpty = 6,              ///< Empty (no-op) node
+  hipGraphNodeTypeWaitEvent = 7,          ///< External event wait node
+  hipGraphNodeTypeEventRecord = 8,        ///< External event record node
+  hipGraphNodeTypeMemcpy1D = 9,           ///< Memcpy 1D node
+  hipGraphNodeTypeMemcpyFromSymbol = 10,  ///< MemcpyFromSymbol node
+  hipGraphNodeTypeMemcpyToSymbol = 11,    ///< MemcpyToSymbol node
+  hipGraphNodeTypeCount
+} hipGraphNodeType;
+
+typedef void (*hipHostFn_t)(void* userData);
+typedef struct hipHostNodeParams {
+  hipHostFn_t fn;
+  void* userData;
+} hipHostNodeParams;
+
+typedef struct hipKernelNodeParams {
+  dim3 blockDim;
+  void** extra;
+  void* func;
+  dim3 gridDim;
+  void** kernelParams;
+  unsigned int sharedMemBytes;
+} hipKernelNodeParams;
+
+typedef struct hipMemsetParams {
+  void* dst;
+  unsigned int elementSize;
+  size_t height;
+  size_t pitch;
+  unsigned int value;
+  size_t width;
+} hipMemsetParams;
+
+enum hipGraphExecUpdateResult {
+  hipGraphExecUpdateSuccess = 0x0,  ///< The update succeeded
+  hipGraphExecUpdateError = 0x1,  ///< The update failed for an unexpected reason which is described
+                                  ///< in the return value of the function
+  hipGraphExecUpdateErrorTopologyChanged = 0x2,  ///< The update failed because the topology changed
+  hipGraphExecUpdateErrorNodeTypeChanged = 0x3,  ///< The update failed because a node type changed
+  hipGraphExecUpdateErrorFunctionChanged =
+      0x4,  ///< The update failed because the function of a kernel node changed
+  hipGraphExecUpdateErrorParametersChanged =
+      0x5,  ///< The update failed because the parameters changed in a way that is not supported
+  hipGraphExecUpdateErrorNotSupported =
+      0x6,  ///< The update failed because something about the node is not supported
+  hipGraphExecUpdateErrorUnsupportedFunctionChange = 0x7
+};
+
+enum hipStreamCaptureMode {
+  hipStreamCaptureModeGlobal = 0,
+  hipStreamCaptureModeThreadLocal,
+  hipStreamCaptureModeRelaxed
+};
+
+enum hipStreamCaptureStatus {
+  hipStreamCaptureStatusNone = 0,    ///< Stream is not capturing
+  hipStreamCaptureStatusActive,      ///< Stream is actively capturing
+  hipStreamCaptureStatusInvalidated  ///< Stream is part of a capture sequence that has been
+                                     ///< invalidated, but not terminated
+};
+
+hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode);
+
+hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph);
+
+// Creates a graph.
+hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags);
+
+// Destroys a graph.
+hipError_t hipGraphDestroy(hipGraph_t graph);
+
+// Destroys an executable graph.
+hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec);
+
+// Creates an executable graph from a graph.
+hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
+                               hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize);
+
+// Launches an executable graph in a stream.
+hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream);
+
+// Creates a kernel execution node and adds it to a graph.
+hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                 const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                 const hipKernelNodeParams* pNodeParams);
+
+// Creates a memcpy node and adds it to a graph.
+hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                 const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                 const hipMemcpy3DParms* pCopyParams);
+
+// Creates a memset node and adds it to a graph.
+hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                 const hipGraphNode_t* pDependencies, size_t numDependencies,
+                                 const hipMemsetParams* pMemsetParams);
+#endif
+// doxygen end graph API
+/**
+ * @}
+ */
 #ifdef __cplusplus
 } /* extern "c" */
 #endif
@@ -4085,6 +4209,11 @@ static inline hipError_t hipUnbindTexture(
     return hipUnbindTexture(&tex);
 }
 
+// doxygen end Texture
+/**
+ * @}
+ */
+
 #endif // __cplusplus
 
 #ifdef __GNUC__

projects/hip/include/hip/hip_runtime_api.h

@@ -280,6 +280,28 @@ typedef enum __HIP_NODISCARD hipError_t {
               ///< that was launched via cooperative launch APIs exceeds the maximum number of
               ///< allowed blocks for the current device
     hipErrorNotSupported = 801,  ///< Produced when the hip API is not supported/implemented
+    hipErrorStreamCaptureUnsupported = 900,  ///< The operation is not permitted when the stream
+                                             ///< is capturing.
+    hipErrorStreamCaptureInvalidated = 901,  ///< The current capture sequence on the stream
+                                             ///< has been invalidated due to a previous error.
+    hipErrorStreamCaptureMerge = 902,  ///< The operation would have resulted in a merge of
+                                       ///< two independent capture sequences.
+    hipErrorStreamCaptureUnmatched = 903,  ///< The capture was not initiated in this stream.
+    hipErrorStreamCaptureUnjoined = 904,  ///< The capture sequence contains a fork that was not
+                                          ///< joined to the primary stream.
+    hipErrorStreamCaptureIsolation = 905,  ///< A dependency would have been created which crosses
+                                           ///< the capture sequence boundary. Only implicit
+                                           ///< in-stream ordering dependencies  are allowed
+                                           ///< to cross the boundary
+    hipErrorStreamCaptureImplicit = 906,  ///< The operation would have resulted in a disallowed
+                                          ///< implicit dependency on a current capture sequence
+                                          ///< from hipStreamLegacy.
+    hipErrorCapturedEvent = 907,  ///< The operation is not permitted on an event which was last
+                                  ///< recorded in a capturing stream.
+    hipErrorStreamCaptureWrongThread = 908,  ///< A stream capture sequence not initiated with
+                                             ///< the hipStreamCaptureModeRelaxed argument to
+                                             ///< hipStreamBeginCapture was passed to 
+                                             ///< hipStreamEndCapture in a different thread.
     hipErrorUnknown = 999,  //< Unknown error.
     // HSA Runtime Error Codes start here.
     hipErrorRuntimeMemory = 1052,  ///< HSA runtime memory call returned error.  Typically not seen