From 6c2759e70de2b90f17e376abac60cc256fbad8da Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Mon, 29 Aug 2016 18:37:57 -0500
Subject: [PATCH] Refactor trace code for hipLaunchKernel.

- Use standard print functions for streams.
- Add HIP_INIT macro, for cases where we want to initialize HIP but not
  log an API (ihipPreKernelLaunch).

Change-Id: If43cf8a363d918bcd3722a2e6a965d4cfa2e03e7


[ROCm/hip commit: 21e5c25225ba6f67b8778e376ad71a08cec5965f]
---
 projects/hip/include/hcc_detail/hip_hcc.h     |  8 ++++++--
 projects/hip/include/hcc_detail/hip_runtime.h | 18 +++++++++---------
 projects/hip/src/hip_hcc.cpp                  |  9 ++++++++-
 3 files changed, 23 insertions(+), 12 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index 449d9bc522..01c36afde4 100644
--- a/projects/hip/include/hcc_detail/hip_hcc.h
+++ b/projects/hip/include/hcc_detail/hip_hcc.h
@@ -160,13 +160,17 @@ class ihipCtx_t;
 #endif
 
 
+// Just initialize the HIP runtime, but don't log any trace information.
+#define HIP_INIT()\
+	std::call_once(hip_initialized, ihipInit);\
+    ihipCtxStackUpdate();
+
 
 // This macro should be called at the beginning of every HIP API.
 // It initialies the hip runtime (exactly once), and
 // generate trace string that can be output to stderr or to ATP file.
 #define HIP_INIT_API(...) \
-	std::call_once(hip_initialized, ihipInit);\
-    ihipCtxStackUpdate();\
+    HIP_INIT()\
     API_TRACE(__VA_ARGS__);
 
 #define ihipLogStatus(hipStatus) \
diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h
index 740e60d0eb..727604d8d8 100644
--- a/projects/hip/include/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hcc_detail/hip_runtime.h
@@ -622,11 +622,12 @@ __device__ static inline void* memset(void* ptr, uint8_t val, size_t size)
 #define HIP_KERNEL_NAME(...) __VA_ARGS__
 
 #ifdef __HCC_CPP__
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp);
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp);
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp);
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp);
-void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
+extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp);
+extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp);
+extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp);
+extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp);
+extern void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream);
+extern void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
 
 // TODO - move to common header file.
 #define KNRM  "\x1B[0m"
@@ -638,10 +639,9 @@ do {\
   grid_launch_parm lp;\
   lp.dynamic_group_mem_bytes = _groupMemBytes; \
   hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp)); \
-    if (HIP_TRACE_API) {\
-        fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:(%d,%d,%d) groupDim:(%d,%d,%d) groupMem:+%d stream=%p\n" KNRM, \
-                #_kernelName, lp.grid_dim.x, lp.grid_dim.y, lp.grid_dim.z, lp.group_dim.x, lp.group_dim.y, lp.group_dim.z, lp.dynamic_group_mem_bytes, (void*)(_stream));\
-    }\
+  if (HIP_TRACE_API) {\
+      ihipPrintKernelLaunch(#_kernelName, &lp, _stream); \
+  }\
   _kernelName (lp, ##__VA_ARGS__);\
   ihipPostLaunchKernel(trueStream, lp);\
 } while(0)
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 09cf46f8a7..340d572bec 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -1314,13 +1314,20 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
     }
 }
 
+void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream) 
+{
+    std::string streamString = ToString(stream);
+    fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:(%d,%d,%d) groupDim:(%d,%d,%d) groupMem:+%d %s\n" KNRM, \
+            kernelName, lp->grid_dim.x, lp->grid_dim.y, lp->grid_dim.z, lp->group_dim.x, lp->group_dim.y, lp->group_dim.z, 
+            lp->dynamic_group_mem_bytes, streamString.c_str());\
+}
 
 // TODO - data-up to data-down:
 // Called just before a kernel is launched from hipLaunchKernel.
 // Allows runtime to track some information about the stream.
 hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp)
 {
-    HIP_INIT_API(stream, grid, block, lp);
+    HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
 #if USE_GRID_LAUNCH_20
     lp->grid_dim.x = grid.x;