SWDEV-240806 - Added API hipLaunchHostFunc and updated graph test for host graph node (#2512)

Change-Id: Idf759064946e503a2b3cc2a38f1a7cb049012688

include/hip/hip_runtime_api.h

@@ -4401,6 +4401,19 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t
                                               size_t numDependencies,
                                               unsigned int flags __dparm(0));
 
+/**
+ * @brief Enqueues a host function call in a stream.
+ *
+ * @param [in] stream - stream to enqueue work to.
+ * @param [in] fn - function to call once operations enqueued preceeding are complete.
+ * @param [in] userData - User-specified data to be passed to the function.
+ * @returns #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorInvalidValue,
+ * #hipErrorNotSupported
+ * @warning : This API is marked as beta, meaning, while this is feature complete,
+ * it is still open to changes and may have outstanding issues.
+ */
+hipError_t hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void* userData);
+
 /**
  * @brief Swaps the stream capture mode of a thread.
  *

tests/src/runtimeApi/graph/hipGraph.cpp

@@ -105,6 +105,25 @@ bool hipWithoutGraphs(float* inputVec_h, float* inputVec_d, double* outputVec_d,
   return true;
 }
 
+typedef struct callBackData {
+  const char* fn_name;
+  double* data;
+} callBackData_t;
+double result_gpu = 0.0;
+void myHostNodeCallback(void* data) {
+  static int iter = 0;
+  iter++;
+  // Check status of GPU after stream operations are done
+  callBackData_t* tmp = (callBackData_t*)(data);
+  // checkCudaErrors(tmp->status);
+  double* result = (double*)(tmp->data);
+  char* function = (char*)(tmp->fn_name);
+  if (iter == GRAPH_LAUNCH_ITERATIONS)
+    printf("[%s] Host callback final reduced sum = %lf\n", function, *result);
+  result_gpu = *result;
+  *result = 0.0;  // reset the result
+}
+
 bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* outputVec_d,
                                  double* result_d, size_t inputSize, size_t numOfBlocks) {
   hipStream_t stream1, stream2, stream3, streamForGraph;
@@ -237,6 +256,16 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d,
   nodeDependencies.clear();
   nodeDependencies.push_back(memcpyNode);
   hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0};
+  hostParams.fn = myHostNodeCallback;
+  callBackData_t hostFnData;
+  hostFnData.data = &result_h;
+  hostFnData.fn_name = "hipGraphsManual";
+  hostParams.userData = &hostFnData;
+
+  HIPCHECK(hipGraphAddHostNode(&hostNode, graph, nodeDependencies.data(), nodeDependencies.size(),
+                               &hostParams));
+
   hipGraphExec_t graphExec;
   hipGraphNode_t* nodes = NULL;
   size_t numNodes = 0;
@@ -266,8 +295,8 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d,
   for (int i = 0; i < inputSize; i++) {
     result_h_cpu += inputVec_h[i];
   }
-  if (result_h_cpu != result_h) {
-    printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_h);
+  if (result_h_cpu != result_gpu) {
+    printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_gpu);
     return false;
   }
   return true;
@@ -304,4 +333,4 @@ int main(int argc, char** argv) {
     failed("Failed during hipGraph with capture\n");
   }
   passed();
-}
+}