From d6aee9ad8ff293e763cba4ffb4e6bed1a7a25162 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Wed, 16 Aug 2017 03:50:04 +0000
Subject: [PATCH] Add some new HIP_TRACE_API options.

---
 hipamd/src/hip_device.cpp     |  2 +-
 hipamd/src/hip_event.cpp      |  6 +++---
 hipamd/src/hip_hcc_internal.h | 10 ++++++----
 hipamd/src/hip_stream.cpp     |  5 +++--
 4 files changed, 13 insertions(+), 10 deletions(-)

diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp
index 1800c9369c..3322690731 100644
--- a/hipamd/src/hip_device.cpp
+++ b/hipamd/src/hip_device.cpp
@@ -153,7 +153,7 @@ hipError_t hipSetDevice(int deviceId)
 
 hipError_t hipDeviceSynchronize(void)
 {
-    HIP_INIT_API();
+    HIP_INIT_SPECIAL_API(TRACE_SYNC);
     return ihipLogStatus(ihipSynchronize());
 }
 
diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp
index 71da9fdc30..3a8f1ab611 100644
--- a/hipamd/src/hip_event.cpp
+++ b/hipamd/src/hip_event.cpp
@@ -127,7 +127,7 @@ hipError_t hipEventCreate(hipEvent_t* event)
 
 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
 {
-    HIP_INIT_API(event, stream);
+    HIP_INIT_SPECIAL_API(TRACE_QUERY, event, stream);
 
     if (event && event->_state != hipEventStatusUnitialized)   {
         stream = ihipSyncAndResolveStream(stream);
@@ -178,7 +178,7 @@ hipError_t hipEventDestroy(hipEvent_t event)
 
 hipError_t hipEventSynchronize(hipEvent_t event)
 {
-    HIP_INIT_API(event);
+    HIP_INIT_SPECIAL_API(TRACE_SYNC, event);
 
     if (event) {
         if (event->_state == hipEventStatusUnitialized) {
@@ -257,7 +257,7 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
 
 hipError_t hipEventQuery(hipEvent_t event)
 {
-    HIP_INIT_API(event);
+    HIP_INIT_SPECIAL_API(TRACE_QUERY, event);
 
     if ((event->_state == hipEventStatusRecording) && (!event->_marker.is_ready())) {
         return ihipLogStatus(hipErrorNotReady);
diff --git a/hipamd/src/hip_hcc_internal.h b/hipamd/src/hip_hcc_internal.h
index 4cb85ffc19..af5e7a121b 100644
--- a/hipamd/src/hip_hcc_internal.h
+++ b/hipamd/src/hip_hcc_internal.h
@@ -191,10 +191,12 @@ extern const char *API_COLOR_END;
 
 //---
 //HIP Trace modes - use with HIP_TRACE_API=...
-#define TRACE_ALL  0 // 0x1
-#define TRACE_KCMD 1 // 0x2, kernel command
-#define TRACE_MCMD 2 // 0x4, memory command
-#define TRACE_MEM  3 // 0x8, memory allocation or deallocation.
+#define TRACE_ALL   0  // 0x01
+#define TRACE_KCMD  1  // 0x02, kernel command
+#define TRACE_MCMD  2  // 0x04, memory command
+#define TRACE_MEM   3  // 0x08, memory allocation or deallocation.
+#define TRACE_SYNC  4  // 0x10, synchronization (host or hipStreamWaitEvent)
+#define TRACE_QUERY 5  // 0x20, hipEventRecord, hipEventQuery, hipStreamQuery
 
 
 //---
diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp
index 40aade28b9..f0e9283201 100644
--- a/hipamd/src/hip_stream.cpp
+++ b/hipamd/src/hip_stream.cpp
@@ -84,7 +84,7 @@ hipError_t hipStreamCreate(hipStream_t *stream)
 
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
 {
-    HIP_INIT_API(stream, event, flags);
+    HIP_INIT_SPECIAL_API(TRACE_SYNC, stream, event, flags);
 
     hipError_t e = hipSuccess;
 
@@ -114,7 +114,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 //---
 hipError_t hipStreamQuery(hipStream_t stream)
 {
-    HIP_INIT_API(stream);
+    HIP_INIT_SPECIAL_API(TRACE_QUERY, stream);
 
     // Use default stream if 0 specified:
     if (stream == hipStreamNull) {
@@ -140,6 +140,7 @@ hipError_t hipStreamQuery(hipStream_t stream)
 hipError_t hipStreamSynchronize(hipStream_t stream)
 {
     HIP_INIT_API(stream);
+    HIP_INIT_SPECIAL_API(TRACE_SYNC, stream);
 
     hipError_t e = hipSuccess;