Merge 'master' into 'amd-master'

Change-Id: I826af9162f47fc666d0de3e7a170c384e93fa2b1

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());
 }
 

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);

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_ALL   0  // 0x01
-#define TRACE_KCMD 1 // 0x2, kernel command
+#define TRACE_KCMD  1  // 0x02, kernel command
-#define TRACE_MCMD 2 // 0x4, memory command
+#define TRACE_MCMD  2  // 0x04, memory command
-#define TRACE_MEM  3 // 0x8, memory allocation or deallocation.
+#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
 
 
 //---

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;