adding activity prof layer

[ROCm/hip commit: b8b1637ef7]

projects/hip/CMakeLists.txt

@@ -150,6 +150,19 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER)
 endif()
 add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
 
+################
+# Detect profiling API
+################
+if (USE_PROF_API EQUAL 1)
+add_definitions(-DUSE_PROF_API=1)
+find_path(PROF_API_HEADER_DIR NAMES prof_protocol.h PATHS ${PROF_API_HEADER_PATH} /opt/rocm/roctracer/include/roctracer NO_DEFAULT_PATH)
+if (NOT PROF_API_HEADER_DIR)
+  MESSAGE("Profiling API header not found, use -DPROF_API_HEADER_PATH=<path to prof_protocol.h header>")
+else ()
+  include_directories ( ${PROF_API_HEADER_DIR} )
+endif ()
+MESSAGE("PROF_API_HEADER_DIR = ${PROF_API_HEADER_DIR}")
+endif ()
 
 #############################
 # Build steps
@@ -204,6 +217,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
         src/grid_launch.cpp
         src/hip_texture.cpp
         src/hip_surface.cpp
+        src/hip_intercept.cpp
         src/env.cpp
         src/program_state.cpp)
 

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -2581,6 +2581,24 @@ hipError_t hipLaunchByPtr(const void* func);
 } /* extern "c" */
 #endif
 
+#include <hip/hip_prof_api.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+/**
+ * Callback/Activity API
+ */
+hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg);
+hipError_t hipRemoveApiCallback(uint32_t id);
+hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg);
+hipError_t hipRemoveActivityCallback(uint32_t id);
+static inline const char* hipApiName(const uint32_t& id) { return hip_api_name(id); }
+const char* hipKernelNameRef(hipFunction_t f);
+#ifdef __cplusplus
+} /* extern "C" */
+#endif
+
 #ifdef __cplusplus
 
 hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,

projects/hip/include/hip/hip_prof_api.h

@@ -0,0 +1,198 @@
+// automatically generated sources
+#ifndef _HIP_PROF_API_H
+#define _HIP_PROF_API_H
+
+#include <atomic>
+#include <mutex>
+
+#include "hip/hip_prof_str.h"
+
+template <typename Record, typename Fun, typename Act>
+class api_callbacks_table_templ {
+ public:
+  typedef std::recursive_mutex mutex_t;
+
+  typedef Record record_t;
+  typedef Fun fun_t;
+  typedef Act act_t;
+
+  // HIP API callbacks table
+  struct hip_cb_table_entry_t {
+    volatile std::atomic<bool> sync;
+    volatile std::atomic<uint32_t> sem;
+    act_t act;
+    void* a_arg;
+    fun_t fun;
+    void* arg;
+  };
+
+  struct hip_cb_table_t {
+    hip_cb_table_entry_t arr[HIP_API_ID_NUMBER];
+  };
+
+  api_callbacks_table_templ() {
+     memset(&callbacks_table_, 0, sizeof(callbacks_table_));
+  }
+
+  bool set_activity(uint32_t id, act_t fun, void* arg) {
+    std::lock_guard<mutex_t> lock(mutex_);
+    bool ret = true;
+    if (id == HIP_API_ID_ANY) {
+      for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_activity(i, fun, arg);
+    } else if (id < HIP_API_ID_NUMBER) {
+      cb_sync(id);
+      callbacks_table_.arr[id].act = fun;
+      callbacks_table_.arr[id].a_arg = arg;
+      cb_release(id);
+    } else {
+      ret = false;
+    }
+    return ret;
+  }
+
+  bool set_callback(uint32_t id, fun_t fun, void* arg) {
+    std::lock_guard<mutex_t> lock(mutex_);
+    bool ret = true;
+    if (id == HIP_API_ID_ANY) {
+      for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_callback(i, fun, arg);
+    } else if (id < HIP_API_ID_NUMBER) {
+      cb_sync(id);
+      callbacks_table_.arr[id].fun = fun;
+      callbacks_table_.arr[id].arg = arg;
+      cb_release(id);
+    } else {
+      ret = false;
+    }
+    return ret;
+  }
+
+  inline hip_cb_table_entry_t& entry(const uint32_t& id) {
+    return callbacks_table_.arr[id];
+  }
+
+  inline void sem_sync(const uint32_t& id) {
+    sem_increment(id);
+    if (entry(id).sync.load() == true) sync_wait(id);
+  }
+
+  inline void sem_release(const uint32_t& id) {
+    sem_decrement(id);
+  }
+
+ private:
+  inline void cb_sync(const uint32_t& id) {
+    entry(id).sync.store(true);
+    while (entry(id).sem.load() != 0) {}
+  }
+
+  inline void cb_release(const uint32_t& id) {
+    entry(id).sync.store(false);
+  }
+
+  inline void sem_increment(const uint32_t& id) {
+    const uint32_t prev = entry(id).sem.fetch_add(1);
+    if (prev == UINT32_MAX) {
+      std::cerr << "sem overflow id = " << id << std::endl << std::flush;
+      abort();
+    }
+  }
+
+  inline void sem_decrement(const uint32_t& id) {
+    const uint32_t prev = entry(id).sem.fetch_sub(1);
+    if (prev == 0) {
+      std::cerr << "sem corrupted id = " << id << std::endl << std::flush;
+      abort();
+    }
+  }
+
+  void sync_wait(const uint32_t& id) {
+    sem_decrement(id);
+    while (entry(id).sync.load() == true) {}
+    sem_increment(id);
+  }
+
+  mutex_t mutex_;
+  hip_cb_table_t callbacks_table_;
+};
+
+
+#if USE_PROF_API
+#include <prof_protocol.h>
+
+static const uint32_t HIP_DOMAIN_ID = ACTIVITY_DOMAIN_HIP_API;
+typedef activity_record_t hip_api_record_t;
+typedef activity_rtapi_callback_t hip_api_callback_t;
+typedef activity_sync_callback_t hip_act_callback_t;
+
+// HIP API callbacks spawner object macro
+#define HIP_CB_SPAWNER_OBJECT(CB_ID) \
+  hip_api_data_t api_data{}; \
+  INIT_CB_ARGS_DATA(CB_ID, api_data); \
+  api_callbacks_spawner_t __api_tracer(HIP_API_ID_##CB_ID, api_data);
+
+typedef api_callbacks_table_templ<hip_api_record_t,
+                                  hip_api_callback_t,
+                                  hip_act_callback_t> api_callbacks_table_t;
+extern api_callbacks_table_t callbacks_table;
+
+class api_callbacks_spawner_t {
+ public:
+  api_callbacks_spawner_t(const hip_api_id_t& cid, hip_api_data_t& api_data) :
+    cid_(cid),
+    api_data_(api_data),
+    record_({})
+  {
+    if (cid_ >= HIP_API_ID_NUMBER) {
+      fprintf(stderr, "HIP %s bad id %d\n", __FUNCTION__, cid_);
+      abort();
+    }
+    callbacks_table.sem_sync(cid_);
+
+    act = entry(cid_).act;
+    a_arg = entry(cid_).a_arg;
+    fun = entry(cid_).fun;
+    arg = entry(cid_).arg;
+
+    api_data_.phase = 0;
+    if (act != NULL) act(cid_, &record_, &api_data_, a_arg);
+    if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, &api_data_, arg);
+  }
+
+  ~api_callbacks_spawner_t() {
+    api_data_.phase = 1;
+    if (act != NULL) act(cid_, &record_, &api_data_, a_arg);
+    if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, &api_data_, arg);
+
+    callbacks_table.sem_release(cid_);
+  }
+
+ private:
+  inline api_callbacks_table_t::hip_cb_table_entry_t& entry(const uint32_t& id) {
+    return callbacks_table.entry(id);
+  }
+
+  const hip_api_id_t cid_;
+  hip_api_data_t& api_data_;
+  hip_api_record_t record_;
+
+  hip_act_callback_t act;
+  void* a_arg;
+  hip_api_callback_t fun;
+  void* arg;
+};
+
+#else
+
+#define HIP_CB_SPAWNER_OBJECT(x) do {} while(0)
+
+class api_callbacks_table_t {
+ public:
+  typedef void* act_t;
+  typedef void* fun_t;
+  bool set_activity(uint32_t id, act_t fun, void* arg) { return true; }
+  bool set_callback(uint32_t id, fun_t fun, void* arg) { return true; }
+};
+
+#endif
+
+#endif  // _HIP_PROF_API_H

projects/hip/src/hip_context.cpp

@@ -40,7 +40,7 @@ void ihipCtxStackUpdate() {
 }
 
 hipError_t hipInit(unsigned int flags) {
-    HIP_INIT_API(flags);
+    HIP_INIT_CB_API(hipInit, flags);
 
     hipError_t e = hipSuccess;
 
@@ -53,7 +53,7 @@ hipError_t hipInit(unsigned int flags) {
 }
 
 hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) {
-    HIP_INIT_API(ctx, flags, device);  // FIXME - review if we want to init
+    HIP_INIT_CB_API(hipCtxCreate, ctx, flags, device);  // FIXME - review if we want to init
     hipError_t e = hipSuccess;
     auto deviceHandle = ihipGetDevice(device);
     {
@@ -71,7 +71,7 @@ hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) {
 }
 
 hipError_t hipDeviceGet(hipDevice_t* device, int deviceId) {
-    HIP_INIT_API(device, deviceId);  // FIXME - review if we want to init
+    HIP_INIT_CB_API(hipDeviceGet, device, deviceId);  // FIXME - review if we want to init
 
     auto deviceHandle = ihipGetDevice(deviceId);
 
@@ -86,7 +86,7 @@ hipError_t hipDeviceGet(hipDevice_t* device, int deviceId) {
 };
 
 hipError_t hipDriverGetVersion(int* driverVersion) {
-    HIP_INIT_API(driverVersion);
+    HIP_INIT_CB_API(hipDriverGetVersion, driverVersion);
     hipError_t e = hipSuccess;
     if (driverVersion) {
         *driverVersion = 4;
@@ -98,7 +98,7 @@ hipError_t hipDriverGetVersion(int* driverVersion) {
 }
 
 hipError_t hipRuntimeGetVersion(int* runtimeVersion) {
-    HIP_INIT_API(runtimeVersion);
+    HIP_INIT_CB_API(hipRuntimeGetVersion, runtimeVersion);
     hipError_t e = hipSuccess;
     if (runtimeVersion) {
         *runtimeVersion = HIP_VERSION_PATCH;
@@ -110,7 +110,7 @@ hipError_t hipRuntimeGetVersion(int* runtimeVersion) {
 }
 
 hipError_t hipCtxDestroy(hipCtx_t ctx) {
-    HIP_INIT_API(ctx);
+    HIP_INIT_CB_API(hipCtxDestroy, ctx);
     hipError_t e = hipSuccess;
     ihipCtx_t* currentCtx = ihipGetTlsDefaultCtx();
     ihipCtx_t* primaryCtx = ((ihipDevice_t*)ctx->getDevice())->_primaryCtx;
@@ -134,7 +134,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) {
 }
 
 hipError_t hipCtxPopCurrent(hipCtx_t* ctx) {
-    HIP_INIT_API(ctx);
+    HIP_INIT_CB_API(hipCtxPopCurrent, ctx);
     hipError_t e = hipSuccess;
     ihipCtx_t* currentCtx = ihipGetTlsDefaultCtx();
     auto deviceHandle = currentCtx->getDevice();
@@ -155,7 +155,7 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) {
 }
 
 hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
-    HIP_INIT_API(ctx);
+    HIP_INIT_CB_API(hipCtxPushCurrent, ctx);
     hipError_t e = hipSuccess;
     if (ctx != NULL) {  // TODO- is this check needed?
         ihipSetTlsDefaultCtx(ctx);
@@ -168,7 +168,7 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
 }
 
 hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
-    HIP_INIT_API(ctx);
+    HIP_INIT_CB_API(hipCtxGetCurrent, ctx);
     hipError_t e = hipSuccess;
     if ((tls_getPrimaryCtx) || tls_ctxStack.empty()) {
         *ctx = ihipGetTlsDefaultCtx();
@@ -179,7 +179,7 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
 }
 
 hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
-    HIP_INIT_API(ctx);
+    HIP_INIT_CB_API(hipCtxSetCurrent, ctx);
     hipError_t e = hipSuccess;
     if (ctx == NULL) {
         tls_ctxStack.pop();
@@ -192,7 +192,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
 }
 
 hipError_t hipCtxGetDevice(hipDevice_t* device) {
-    HIP_INIT_API(device);
+    HIP_INIT_CB_API(hipCtxGetDevice, device);
     hipError_t e = hipSuccess;
 
     ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
@@ -208,7 +208,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) {
 }
 
 hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
-    HIP_INIT_API(apiVersion);
+    HIP_INIT_CB_API(hipCtxGetApiVersion, apiVersion);
 
     if (apiVersion) {
         *apiVersion = 4;
@@ -218,7 +218,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
 }
 
 hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
-    HIP_INIT_API(cacheConfig);
+    HIP_INIT_CB_API(hipCtxGetCacheConfig, cacheConfig);
 
     *cacheConfig = hipFuncCachePreferNone;
 
@@ -226,7 +226,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
 }
 
 hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
-    HIP_INIT_API(cacheConfig);
+    HIP_INIT_CB_API(hipCtxSetCacheConfig, cacheConfig);
 
     // Nop, AMD does not support variable cache configs.
 
@@ -234,7 +234,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
 }
 
 hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
-    HIP_INIT_API(config);
+    HIP_INIT_CB_API(hipCtxSetSharedMemConfig, config);
 
     // Nop, AMD does not support variable shared mem configs.
 
@@ -242,7 +242,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
 }
 
 hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
-    HIP_INIT_API(pConfig);
+    HIP_INIT_CB_API(hipCtxGetSharedMemConfig, pConfig);
 
     *pConfig = hipSharedMemBankSizeFourByte;
 
@@ -250,12 +250,12 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
 }
 
 hipError_t hipCtxSynchronize(void) {
-    HIP_INIT_API(1);
+    HIP_INIT_CB_API(hipCtxSynchronize, 1);
     return ihipLogStatus(ihipSynchronize());  // TODP Shall check validity of ctx?
 }
 
 hipError_t hipCtxGetFlags(unsigned int* flags) {
-    HIP_INIT_API(flags);
+    HIP_INIT_CB_API(hipCtxGetFlags, flags);
     hipError_t e = hipSuccess;
     ihipCtx_t* tempCtx;
     tempCtx = ihipGetTlsDefaultCtx();
@@ -264,7 +264,7 @@ hipError_t hipCtxGetFlags(unsigned int* flags) {
 }
 
 hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active) {
-    HIP_INIT_API(dev, flags, active);
+    HIP_INIT_CB_API(hipDevicePrimaryCtxGetState, dev, flags, active);
     hipError_t e = hipSuccess;
     auto deviceHandle = ihipGetDevice(dev);
 
@@ -286,7 +286,7 @@ hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int
 }
 
 hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) {
-    HIP_INIT_API(dev);
+    HIP_INIT_CB_API(hipDevicePrimaryCtxRelease, dev);
     hipError_t e = hipSuccess;
     auto deviceHandle = ihipGetDevice(dev);
 
@@ -297,7 +297,7 @@ hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) {
 }
 
 hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
-    HIP_INIT_API(pctx, dev);
+    HIP_INIT_CB_API(hipDevicePrimaryCtxRetain, pctx, dev);
     hipError_t e = hipSuccess;
     auto deviceHandle = ihipGetDevice(dev);
 
@@ -309,7 +309,7 @@ hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
 }
 
 hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
-    HIP_INIT_API(dev);
+    HIP_INIT_CB_API(hipDevicePrimaryCtxReset, dev);
     hipError_t e = hipSuccess;
     auto deviceHandle = ihipGetDevice(dev);
 
@@ -322,7 +322,7 @@ hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
 }
 
 hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) {
-    HIP_INIT_API(dev, flags);
+    HIP_INIT_CB_API(hipDevicePrimaryCtxSetFlags, dev, flags);
     hipError_t e = hipSuccess;
     auto deviceHandle = ihipGetDevice(dev);
 

projects/hip/src/hip_device.cpp

@@ -30,7 +30,7 @@ THE SOFTWARE.
 //-------------------------------------------------------------------------------------------------
 // TODO - does this initialize HIP runtime?
 hipError_t hipGetDevice(int* deviceId) {
-    HIP_INIT_API(deviceId);
+    HIP_INIT_CB_API(hipGetDevice, deviceId);
 
     hipError_t e = hipSuccess;
 
@@ -69,12 +69,12 @@ hipError_t ihipGetDeviceCount(int* count) {
 }
 
 hipError_t hipGetDeviceCount(int* count) {
-    HIP_INIT_API(count);
+    HIP_INIT_CB_API(hipGetDeviceCount, count);
     return ihipLogStatus(ihipGetDeviceCount(count));
 }
 
 hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) {
-    HIP_INIT_API(cacheConfig);
+    HIP_INIT_CB_API(hipDeviceSetCacheConfig, cacheConfig);
 
     // Nop, AMD does not support variable cache configs.
 
@@ -82,7 +82,7 @@ hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) {
 }
 
 hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig) {
-    HIP_INIT_API(cacheConfig);
+    HIP_INIT_CB_API(hipDeviceGetCacheConfig, cacheConfig);
 
     if (cacheConfig == nullptr) {
         return ihipLogStatus(hipErrorInvalidValue);
@@ -94,7 +94,7 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig) {
 }
 
 hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) {
-    HIP_INIT_API(pValue, limit);
+    HIP_INIT_CB_API(hipDeviceGetLimit, pValue, limit);
     if (pValue == nullptr) {
         return ihipLogStatus(hipErrorInvalidValue);
     }
@@ -107,7 +107,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) {
 }
 
 hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) {
-    HIP_INIT_API(cacheConfig);
+    HIP_INIT_CB_API(hipFuncSetCacheConfig, cacheConfig);
 
     // Nop, AMD does not support variable cache configs.
 
@@ -115,7 +115,7 @@ hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) {
 }
 
 hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config) {
-    HIP_INIT_API(config);
+    HIP_INIT_CB_API(hipDeviceSetSharedMemConfig, config);
 
     // Nop, AMD does not support variable shared mem configs.
 
@@ -123,7 +123,7 @@ hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config) {
 }
 
 hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig) {
-    HIP_INIT_API(pConfig);
+    HIP_INIT_CB_API(hipDeviceGetSharedMemConfig, pConfig);
 
     *pConfig = hipSharedMemBankSizeFourByte;
 
@@ -131,7 +131,7 @@ hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig) {
 }
 
 hipError_t hipSetDevice(int deviceId) {
-    HIP_INIT_API(deviceId);
+    HIP_INIT_CB_API(hipSetDevice, deviceId);
     if ((deviceId < 0) || (deviceId >= g_deviceCnt)) {
         return ihipLogStatus(hipErrorInvalidDevice);
     } else {
@@ -142,12 +142,12 @@ hipError_t hipSetDevice(int deviceId) {
 }
 
 hipError_t hipDeviceSynchronize(void) {
-    HIP_INIT_SPECIAL_API(TRACE_SYNC);
+    HIP_INIT_SPECIAL_CB_API(hipDeviceSynchronize, TRACE_SYNC);
     return ihipLogStatus(ihipSynchronize());
 }
 
 hipError_t hipDeviceReset(void) {
-    HIP_INIT_API();
+    HIP_INIT_CB_API(hipDeviceReset, );
 
     auto* ctx = ihipGetTlsDefaultCtx();
 
@@ -287,7 +287,7 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device
 }
 
 hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) {
-    HIP_INIT_API(pi, attr, device);
+    HIP_INIT_CB_API(hipDeviceGetAttribute, pi, attr, device);
     if ((device < 0) || (device >= g_deviceCnt)) {
         return ihipLogStatus(hipErrorInvalidDevice);
     }
@@ -314,7 +314,7 @@ hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device) {
 }
 
 hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) {
-    HIP_INIT_API(props, device);
+    HIP_INIT_CB_API(hipGetDeviceProperties, props, device);
     if ((device < 0) || (device >= g_deviceCnt)) {
         return ihipLogStatus(hipErrorInvalidDevice);
     }
@@ -322,7 +322,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) {
 }
 
 hipError_t hipSetDeviceFlags(unsigned int flags) {
-    HIP_INIT_API(flags);
+    HIP_INIT_CB_API(hipSetDeviceFlags, flags);
 
     hipError_t e = hipSuccess;
 
@@ -367,7 +367,7 @@ hipError_t hipSetDeviceFlags(unsigned int flags) {
 };
 
 hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device) {
-    HIP_INIT_API(major, minor, device);
+    HIP_INIT_CB_API(hipDeviceComputeCapability, major, minor, device);
     hipError_t e = hipSuccess;
     if ((device < 0) || (device >= g_deviceCnt)) {
         e = hipErrorInvalidDevice;
@@ -380,7 +380,7 @@ hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device
 
 hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device) {
     // Cast to void* here to avoid printing garbage in debug modes.
-    HIP_INIT_API((void*)name, len, device);
+    HIP_INIT_CB_API(hipDeviceGetName, (void*)name, len, device);
     hipError_t e = hipSuccess;
     if ((device < 0) || (device >= g_deviceCnt)) {
         e = hipErrorInvalidDevice;
@@ -394,7 +394,7 @@ hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device) {
 
 hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) {
     // Cast to void* here to avoid printing garbage in debug modes.
-    HIP_INIT_API((void*)pciBusId, len, device);
+    HIP_INIT_CB_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device);
     hipError_t e = hipErrorInvalidValue;
     if ((device < 0) || (device >= g_deviceCnt)) {
         e = hipErrorInvalidDevice;
@@ -413,7 +413,7 @@ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) {
 }
 
 hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device) {
-    HIP_INIT_API(bytes, device);
+    HIP_INIT_CB_API(hipDeviceTotalMem, bytes, device);
     hipError_t e = hipSuccess;
     if ((device < 0) || (device >= g_deviceCnt)) {
         e = hipErrorInvalidDevice;
@@ -425,7 +425,7 @@ hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device) {
 }
 
 hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId) {
-    HIP_INIT_API(device, pciBusId);
+    HIP_INIT_CB_API(hipDeviceGetByPCIBusId, device, pciBusId);
     hipDeviceProp_t tempProp;
     int deviceCount = 0;
     hipError_t e = hipErrorInvalidValue;
@@ -451,7 +451,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId) {
 }
 
 hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop) {
-    HIP_INIT_API(device, prop);
+    HIP_INIT_CB_API(hipChooseDevice, device, prop);
     hipDeviceProp_t tempProp;
     hipError_t e = hipSuccess;
     if ((device == NULL) || (prop == NULL)) {

projects/hip/src/hip_error.cpp

@@ -30,7 +30,7 @@ THE SOFTWARE.
 //---
 
 hipError_t hipGetLastError() {
-    HIP_INIT_API();
+    HIP_INIT_CB_API(hipGetLastError);
 
     // Return last error, but then reset the state:
     hipError_t e = ihipLogStatus(tls_lastHipError);
@@ -39,7 +39,7 @@ hipError_t hipGetLastError() {
 }
 
 hipError_t hipPeekAtLastError() {
-    HIP_INIT_API();
+    HIP_INIT_CB_API(hipPeekAtLastError);
 
     // peek at last error, but don't reset it.
     return ihipLogStatus(tls_lastHipError);

projects/hip/src/hip_event.cpp

@@ -95,20 +95,20 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) {
 }
 
 hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
-    HIP_INIT_API(event, flags);
+    HIP_INIT_CB_API(hipEventCreateWithFlags, event, flags);
 
     return ihipLogStatus(ihipEventCreate(event, flags));
 }
 
 hipError_t hipEventCreate(hipEvent_t* event) {
-    HIP_INIT_API(event);
+    HIP_INIT_CB_API(hipEventCreate, event);
 
     return ihipLogStatus(ihipEventCreate(event, 0));
 }
 
 
 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
-    HIP_INIT_SPECIAL_API(TRACE_SYNC, event, stream);
+    HIP_INIT_SPECIAL_CB_API(hipEventRecord, TRACE_SYNC, event, stream);
 
     auto ecd = event->locked_copyCrit();
 
@@ -153,7 +153,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
 
 
 hipError_t hipEventDestroy(hipEvent_t event) {
-    HIP_INIT_API(event);
+    HIP_INIT_CB_API(hipEventDestroy, event);
 
     if (event) {
         delete event;
@@ -165,7 +165,7 @@ hipError_t hipEventDestroy(hipEvent_t event) {
 }
 
 hipError_t hipEventSynchronize(hipEvent_t event) {
-    HIP_INIT_SPECIAL_API(TRACE_SYNC, event);
+    HIP_INIT_SPECIAL_CB_API(hipEventSynchronize, TRACE_SYNC, event);
 
     if (!(event->_flags & hipEventReleaseToSystem)) {
         tprintf(DB_WARN,
@@ -198,7 +198,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
 }
 
 hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
-    HIP_INIT_API(ms, start, stop);
+    HIP_INIT_CB_API(hipEventElapsedTime, ms, start, stop);
 
     hipError_t status = hipSuccess;
 
@@ -255,7 +255,7 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
 }
 
 hipError_t hipEventQuery(hipEvent_t event) {
-    HIP_INIT_SPECIAL_API(TRACE_QUERY, event);
+    HIP_INIT_SPECIAL_CB_API(hipEventQuery, TRACE_QUERY, event);
 
     if (!(event->_flags & hipEventReleaseToSystem)) {
         tprintf(DB_WARN,

projects/hip/src/hip_hcc.cpp

@@ -2288,7 +2288,7 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
 //-------------------------------------------------------------------------------------------------
 // Profiler, really these should live elsewhere:
 hipError_t hipProfilerStart() {
-    HIP_INIT_API();
+    HIP_INIT_CB_API(hipProfilerStart);
 #if COMPILE_HIP_ATP_MARKER
     amdtResumeProfiling(AMDT_ALL_PROFILING);
 #endif
@@ -2298,7 +2298,7 @@ hipError_t hipProfilerStart() {
 
 
 hipError_t hipProfilerStop() {
-    HIP_INIT_API();
+    HIP_INIT_CB_API(hipProfilerStop);
 #if COMPILE_HIP_ATP_MARKER
     amdtStopProfiling(AMDT_ALL_PROFILING);
 #endif

projects/hip/src/hip_hcc_internal.h

@@ -301,6 +301,11 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr);
     HIP_INIT()                                                                                     \
     API_TRACE(0, __VA_ARGS__);
 
+#define HIP_INIT_CB_API(cid, ...)                                                                  \
+    HIP_INIT()                                                                                     \
+    API_TRACE(0, __VA_ARGS__);                                                                     \
+    HIP_CB_SPAWNER_OBJECT(cid);
+
 
 // Like above, but will trace with a specified "special" bit.
 // Replace HIP_INIT_API with this call inside HIP APIs that launch work on the GPU:
@@ -309,6 +314,11 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr);
     HIP_INIT()                                                                                     \
     API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__);
 
+#define HIP_INIT_SPECIAL_CB_API(cid, tbit, ...)                                                    \
+    HIP_INIT()                                                                                     \
+    API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__);                                         \
+    HIP_CB_SPAWNER_OBJECT(cid);
+
 
 // This macro should be called at the end of every HIP API, and only at the end of top-level hip
 // APIS (not internal hip) It has dual function: logs the last error returned for use by

projects/hip/src/hip_intercept.cpp

@@ -0,0 +1,49 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+#include "hip/hip_prof_api.h"
+
+// HIP API callback/activity
+
+api_callbacks_table_t callbacks_table;
+
+extern std::string& FunctionSymbol(hipFunction_t f);
+const char* hipKernelNameRef(const hipFunction_t f) { return FunctionSymbol(f).c_str(); }
+
+hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg) {
+  return callbacks_table.set_callback(id, reinterpret_cast<api_callbacks_table_t::fun_t>(fun), arg) ?
+    hipSuccess : hipErrorInvalidValue;
+}
+
+hipError_t hipRemoveApiCallback(uint32_t id) {
+  return callbacks_table.set_callback(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue;
+}
+
+hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg) {
+  return callbacks_table.set_activity(id, reinterpret_cast<api_callbacks_table_t::act_t>(fun), arg) ?
+    hipSuccess : hipErrorInvalidValue;
+}
+
+hipError_t hipRemoveActivityCallback(uint32_t id) {
+  return callbacks_table.set_activity(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue;
+}

projects/hip/src/hip_memory.cpp

@@ -159,7 +159,7 @@ void* allocAndSharePtr(const char* msg, size_t sizeBytes, ihipCtx_t* ctx, bool s
 // TODO - add more info here when available.
 //
 hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) {
-    HIP_INIT_API(attributes, ptr);
+    HIP_INIT_CB_API(hipPointerGetAttributes, attributes, ptr);
 
     hipError_t e = hipSuccess;
     if ((attributes == nullptr) || (ptr == nullptr)) {
@@ -206,7 +206,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void
 
 
 hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) {
-    HIP_INIT_API(devicePointer, hostPointer, flags);
+    HIP_INIT_CB_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags);
 
     hipError_t e = hipSuccess;
 
@@ -237,7 +237,7 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi
 
 
 hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, sizeBytes);
+    HIP_INIT_SPECIAL_CB_API(hipMalloc, (TRACE_MEM), ptr, sizeBytes);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
 
@@ -266,7 +266,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
 
 
 hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, sizeBytes, flags);
+    HIP_INIT_SPECIAL_CB_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
 
@@ -398,7 +398,7 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh
 
 // width in bytes
 hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, pitch, width, height);
+    HIP_INIT_SPECIAL_CB_API(hipMallocPitch, (TRACE_MEM), ptr, pitch, width, height);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
 
@@ -409,7 +409,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
 }
 
 hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
-    HIP_INIT_API(pitchedDevPtr, &extent);
+    HIP_INIT_CB_API(hipMalloc3D, pitchedDevPtr, &extent);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
 
@@ -444,7 +444,7 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc,
                                    hsa_ext_image_channel_type_t* channelType);
 
 hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), array, pAllocateArray);
+    HIP_INIT_SPECIAL_CB_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
     if (pAllocateArray->width > 0) {
@@ -554,7 +554,7 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat
 
 hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width,
                           size_t height, unsigned int flags) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), array, desc, width, height, flags);
+    HIP_INIT_SPECIAL_CB_API(hipMallocArray, (TRACE_MEM), array, desc, width, height, flags);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
     if (width > 0) {
@@ -635,7 +635,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si
 }
 
 hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), array, pAllocateArray);
+    HIP_INIT_SPECIAL_CB_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray);
     hipError_t hip_status = hipSuccess;
 
     auto ctx = ihipGetTlsDefaultCtx();
@@ -762,7 +762,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
 
 
 
-    HIP_INIT_API(array, desc, &extent, flags);
+    HIP_INIT_CB_API(hipMalloc3DArray, array, desc, &extent, flags);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
 
@@ -846,7 +846,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
 }
 
 hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
-    HIP_INIT_API(flagsPtr, hostPtr);
+    HIP_INIT_CB_API(hipHostGetFlags, flagsPtr, hostPtr);
 
     hipError_t hip_status = hipSuccess;
 
@@ -874,7 +874,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
 
 // TODO - need to fix several issues here related to P2P access, host memory fallback.
 hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) {
-    HIP_INIT_API(hostPtr, sizeBytes, flags);
+    HIP_INIT_CB_API(hipHostRegister, hostPtr, sizeBytes, flags);
 
     hipError_t hip_status = hipSuccess;
 
@@ -931,7 +931,7 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
 }
 
 hipError_t hipHostUnregister(void* hostPtr) {
-    HIP_INIT_API(hostPtr);
+    HIP_INIT_CB_API(hipHostUnregister, hostPtr);
     auto ctx = ihipGetTlsDefaultCtx();
     hipError_t hip_status = hipSuccess;
     if (hostPtr == NULL) {
@@ -966,7 +966,7 @@ inline hipDeviceptr_t agent_address_for_symbol(const char* symbolName) {
 
 hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count, size_t offset,
                              hipMemcpyKind kind) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyToSymbol, (TRACE_MCMD), symbolName, src, count, offset, kind);
 
     if (symbolName == nullptr) {
         return ihipLogStatus(hipErrorInvalidSymbol);
@@ -999,7 +999,7 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou
 
 hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, size_t offset,
                                hipMemcpyKind kind) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, dst, count, offset, kind);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyFromSymbol, (TRACE_MCMD), symbolName, dst, count, offset, kind);
 
     if (symbolName == nullptr) {
         return ihipLogStatus(hipErrorInvalidSymbol);
@@ -1031,7 +1031,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count,
 
 hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_t count,
                                   size_t offset, hipMemcpyKind kind, hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyToSymbolAsync, (TRACE_MCMD), symbolName, src, count, offset, kind, stream);
 
     if (symbolName == nullptr) {
         return ihipLogStatus(hipErrorInvalidSymbol);
@@ -1066,7 +1066,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_
 
 hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count, size_t offset,
                                     hipMemcpyKind kind, hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, dst, count, offset, kind, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyFromSymbolAsync, (TRACE_MCMD), symbolName, dst, count, offset, kind, stream);
 
     if (symbolName == nullptr) {
         return ihipLogStatus(hipErrorInvalidSymbol);
@@ -1101,7 +1101,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
 
 //---
 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind);
 
     hipError_t e = hipSuccess;
 
@@ -1128,7 +1128,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
 
 
 hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1147,7 +1147,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
 
 
 hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1166,7 +1166,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
 
 
 hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1203,13 +1203,13 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
 
 hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
                           hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyAsync, (TRACE_MCMD), dst, src, sizeBytes, kind, stream);
 
     return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream));
 }
 
 hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyHtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream);
 
     return ihipLogStatus(
         hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyHostToDevice, stream));
@@ -1217,14 +1217,14 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, h
 
 hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes,
                               hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyDtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream);
 
     return ihipLogStatus(
         hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream));
 }
 
 hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyDtoHAsync, (TRACE_MCMD), dst, src, sizeBytes, stream);
 
     return ihipLogStatus(
         hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream));
@@ -1232,7 +1232,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
 
 hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
                               size_t spitch, size_t width, size_t height, hipMemcpyKind kind) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpy2DToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1284,7 +1284,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
 
 hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
                             size_t count, hipMemcpyKind kind) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, count, kind);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, count, kind);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1303,7 +1303,7 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const
 
 hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
                               size_t count, hipMemcpyKind kind) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyFromArray, (TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1321,7 +1321,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs
 }
 
 hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dstArray, dstOffset, srcHost, count);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyHtoA, (TRACE_MCMD), dstArray, dstOffset, srcHost, count);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1339,7 +1339,7 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo
 }
 
 hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, srcOffset, count);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyAtoH, (TRACE_MCMD), dst, srcArray, srcOffset, count);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -1358,7 +1358,7 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
 }
 
 hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), p);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpy3D, (TRACE_MCMD), p);
     hipError_t e = hipSuccess;
     if (p) {
         size_t byteSize;
@@ -1626,7 +1626,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
 
 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                        size_t height, hipMemcpyKind kind) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpy2D, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
     hipError_t e = hipSuccess;
     e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind);
     return ihipLogStatus(e);
@@ -1634,7 +1634,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
 
 hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                             size_t height, hipMemcpyKind kind, hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
     if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue);
     hipError_t e = hipSuccess;
     int isLocked = 0;
@@ -1673,7 +1673,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
 }
 
 hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy);
+    HIP_INIT_SPECIAL_CB_API(hipMemcpyParam2D, (TRACE_MCMD), pCopy);
     hipError_t e = hipSuccess;
     if (pCopy == nullptr) {
         e = hipErrorInvalidValue;
@@ -1685,7 +1685,7 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
 
 // TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
 hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemsetAsync, (TRACE_MCMD), dst, value, sizeBytes, stream);
 
     hipError_t e = hipSuccess;
 
@@ -1697,7 +1697,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
 };
 
 hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
+    HIP_INIT_SPECIAL_CB_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes);
 
     hipError_t e = hipSuccess;
 
@@ -1713,7 +1713,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
 }
 
 hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, pitch, value, width, height);
+    HIP_INIT_SPECIAL_CB_API(hipMemset2D, (TRACE_MCMD), dst, pitch, value, width, height);
 
     hipError_t e = hipSuccess;
 
@@ -1732,7 +1732,7 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
 
 hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream )
 {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, pitch, value, width, height, stream);
+    HIP_INIT_SPECIAL_CB_API(hipMemset2DAsync, (TRACE_MCMD), dst, pitch, value, width, height, stream);
 
     hipError_t e = hipSuccess;
 
@@ -1749,7 +1749,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, si
 };
 
 hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
+    HIP_INIT_SPECIAL_CB_API(hipMemsetD8, (TRACE_MCMD), dst, value, sizeBytes);
 
     hipError_t e = hipSuccess;
 
@@ -1766,7 +1766,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
 
 hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int  value, hipExtent extent )
 {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent);
+    HIP_INIT_SPECIAL_CB_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent);
     hipError_t e = hipSuccess;
 
     hipStream_t stream = hipStreamNull;
@@ -1785,7 +1785,7 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int  value, hipExtent extent
 
 hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int  value, hipExtent extent ,hipStream_t stream )
 {
-    HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent);
+    HIP_INIT_SPECIAL_CB_API(hipMemset3DAsync, (TRACE_MCMD), &pitchedDevPtr, value, &extent);
     hipError_t e = hipSuccess;
 
     // TODO - call an ihip memset so HIP_TRACE is correct.
@@ -1801,7 +1801,7 @@ hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int  value, hipExtent e
 }
 
 hipError_t hipMemGetInfo(size_t* free, size_t* total) {
-    HIP_INIT_API(free, total);
+    HIP_INIT_CB_API(hipMemGetInfo, free, total);
 
     hipError_t e = hipSuccess;
 
@@ -1835,7 +1835,7 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) {
 }
 
 hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) {
-    HIP_INIT_API(ptr, size);
+    HIP_INIT_CB_API(hipMemPtrGetInfo, ptr, size);
 
     hipError_t e = hipSuccess;
 
@@ -1860,7 +1860,7 @@ hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) {
 
 
 hipError_t hipFree(void* ptr) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
+    HIP_INIT_SPECIAL_CB_API(hipFree, (TRACE_MEM), ptr);
 
     hipError_t hipStatus = hipErrorInvalidDevicePointer;
 
@@ -1892,7 +1892,7 @@ hipError_t hipFree(void* ptr) {
 
 
 hipError_t hipHostFree(void* ptr) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
+    HIP_INIT_SPECIAL_CB_API(hipHostFree, (TRACE_MEM), ptr);
 
     // Synchronize to ensure all work has finished.
     ihipGetTlsDefaultCtx()->locked_waitAllStreams();  // ignores non-blocking streams, this waits
@@ -1927,7 +1927,7 @@ hipError_t hipHostFree(void* ptr) {
 hipError_t hipFreeHost(void* ptr) { return hipHostFree(ptr); }
 
 hipError_t hipFreeArray(hipArray* array) {
-    HIP_INIT_SPECIAL_API((TRACE_MEM), array);
+    HIP_INIT_SPECIAL_CB_API(hipFreeArray, (TRACE_MEM), array);
 
     hipError_t hipStatus = hipErrorInvalidDevicePointer;
 
@@ -1955,7 +1955,7 @@ hipError_t hipFreeArray(hipArray* array) {
 }
 
 hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) {
-    HIP_INIT_API(pbase, psize, dptr);
+    HIP_INIT_CB_API(hipMemGetAddressRange, pbase, psize, dptr);
     hipError_t hipStatus = hipSuccess;
     hc::accelerator acc;
 #if (__hcc_workweek__ >= 17332)
@@ -1976,7 +1976,7 @@ hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDevice
 // TODO: IPC implementaiton:
 
 hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
-    HIP_INIT_API(handle, devPtr);
+    HIP_INIT_CB_API(hipIpcGetMemHandle, handle, devPtr);
     hipError_t hipStatus = hipSuccess;
     // Get the size of allocated pointer
     size_t psize = 0u;
@@ -2012,7 +2012,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
 }
 
 hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags) {
-    HIP_INIT_API(devPtr, &handle, flags);
+    HIP_INIT_CB_API(hipIpcOpenMemHandle, devPtr, &handle, flags);
     hipError_t hipStatus = hipSuccess;
     if (devPtr == NULL) {
         hipStatus = hipErrorInvalidValue;
@@ -2042,7 +2042,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
 }
 
 hipError_t hipIpcCloseMemHandle(void* devPtr) {
-    HIP_INIT_API(devPtr);
+    HIP_INIT_CB_API(hipIpcCloseMemHandle, devPtr);
     hipError_t hipStatus = hipSuccess;
     if (devPtr == NULL) {
         hipStatus = hipErrorInvalidValue;

projects/hip/src/hip_module.cpp

@@ -96,7 +96,7 @@ string ToString(hipFunction_t v) {
     }
 
 hipError_t hipModuleUnload(hipModule_t hmod) {
-    HIP_INIT_API(hmod);
+    HIP_INIT_CB_API(hipModuleUnload, hmod);
 
     // TODO - improve this synchronization so it is thread-safe.
     // Currently we want for all inflight activity to complete, but don't prevent another
@@ -230,7 +230,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr
                                  uint32_t gridDimZ, uint32_t blockDimX, uint32_t blockDimY,
                                  uint32_t blockDimZ, uint32_t sharedMemBytes, hipStream_t hStream,
                                  void** kernelParams, void** extra) {
-    HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes,
+    HIP_INIT_CB_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes,
                  hStream, kernelParams, extra);
     return ihipLogStatus(ihipModuleLaunchKernel(
         f, blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ, blockDimX, blockDimY,
@@ -460,13 +460,13 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch
 }
 
 hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) {
-    HIP_INIT_API(hfunc, hmod, name);
+    HIP_INIT_CB_API(hipModuleGetFunction, hfunc, hmod, name);
     return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name));
 }
 
 hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
                               const char* name) {
-    HIP_INIT_API(dptr, bytes, hmod, name);
+    HIP_INIT_CB_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
 
     if (!dptr || !bytes) return ihipLogStatus(hipErrorInvalidValue);
 
@@ -558,12 +558,12 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
 }
 
 hipError_t hipModuleLoadData(hipModule_t* module, const void* image) {
-    HIP_INIT_API(module, image);
+    HIP_INIT_CB_API(hipModuleLoadData, module, image);
     return ihipLogStatus(ihipModuleLoadData(module,image));
 }
 
 hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
-    HIP_INIT_API(module, fname);
+    HIP_INIT_CB_API(hipModuleLoad, module, fname);
 
     if (!fname) return ihipLogStatus(hipErrorInvalidValue);
 
@@ -578,12 +578,12 @@ hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
 
 hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions,
                                hipJitOption* options, void** optionValues) {
-    HIP_INIT_API(module, image, numOptions, options, optionValues);
+    HIP_INIT_CB_API(hipModuleLoadDataEx, module, image, numOptions, options, optionValues);
     return ihipLogStatus(ihipModuleLoadData(module, image));
 }
 
 hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) {
-    HIP_INIT_API(texRef, hmod, name);
+    HIP_INIT_CB_API(hipModuleGetTexRef, texRef, hmod, name);
 
     hipError_t ret = hipErrorNotFound;
     if (!texRef) return ihipLogStatus(hipErrorInvalidValue);

projects/hip/src/hip_peer.cpp

@@ -175,21 +175,21 @@ hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hi
 //=============================================================================
 
 hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) {
-    HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId);
+    HIP_INIT_CB_API(hipDeviceCanAccessPeer, canAccessPeer, deviceId, peerDeviceId);
     return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId),
                                                  ihipGetPrimaryCtx(peerDeviceId)));
 }
 
 
 hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) {
-    HIP_INIT_API(peerDeviceId);
+    HIP_INIT_CB_API(hipDeviceDisablePeerAccess, peerDeviceId);
 
     return ihipLogStatus(ihipDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId)));
 }
 
 
 hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) {
-    HIP_INIT_API(peerDeviceId, flags);
+    HIP_INIT_CB_API(hipDeviceEnablePeerAccess, peerDeviceId, flags);
 
     return ihipLogStatus(ihipEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags));
 }
@@ -197,7 +197,7 @@ hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) {
 
 hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice,
                          size_t sizeBytes) {
-    HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes);
+    HIP_INIT_CB_API(hipMemcpyPeer, dst, dstDevice, src, srcDevice, sizeBytes);
     return ihipLogStatus(hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src,
                                        ihipGetPrimaryCtx(srcDevice), sizeBytes));
 }
@@ -205,18 +205,18 @@ hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevic
 
 hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice,
                               size_t sizeBytes, hipStream_t stream) {
-    HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
+    HIP_INIT_CB_API(hipMemcpyPeerAsync, dst, dstDevice, src, srcDevice, sizeBytes, stream);
     return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream));
 }
 
 hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
-    HIP_INIT_API(peerCtx, flags);
+    HIP_INIT_CB_API(hipCtxEnablePeerAccess, peerCtx, flags);
 
     return ihipLogStatus(ihipEnablePeerAccess(peerCtx, flags));
 }
 
 hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) {
-    HIP_INIT_API(peerCtx);
+    HIP_INIT_CB_API(hipCtxDisablePeerAccess, peerCtx);
 
     return ihipLogStatus(ihipDisablePeerAccess(peerCtx));
 }

projects/hip/src/hip_stream.cpp

@@ -90,14 +90,14 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
 
 //---
 hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) {
-    HIP_INIT_API(stream, flags);
+    HIP_INIT_CB_API(hipStreamCreateWithFlags, stream, flags);
 
     return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
 }
 
 //---
 hipError_t hipStreamCreate(hipStream_t* stream) {
-    HIP_INIT_API(stream);
+    HIP_INIT_CB_API(hipStreamCreate, stream);
 
     return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault, priority_normal));
 }
@@ -121,7 +121,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio
 }
 
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
-    HIP_INIT_SPECIAL_API(TRACE_SYNC, stream, event, flags);
+    HIP_INIT_SPECIAL_CB_API(hipStreamWaitEvent, TRACE_SYNC, stream, event, flags);
 
     hipError_t e = hipSuccess;
 
@@ -152,7 +152,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 
 //---
 hipError_t hipStreamQuery(hipStream_t stream) {
-    HIP_INIT_SPECIAL_API(TRACE_QUERY, stream);
+    HIP_INIT_SPECIAL_CB_API(hipStreamQuery, TRACE_QUERY, stream);
 
     // Use default stream if 0 specified:
     if (stream == hipStreamNull) {
@@ -175,7 +175,7 @@ hipError_t hipStreamQuery(hipStream_t stream) {
 
 //---
 hipError_t hipStreamSynchronize(hipStream_t stream) {
-    HIP_INIT_SPECIAL_API(TRACE_SYNC, stream);
+    HIP_INIT_SPECIAL_CB_API(hipStreamSynchronize, TRACE_SYNC, stream);
 
     return ihipLogStatus(ihipStreamSynchronize(stream));
 }
@@ -186,7 +186,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) {
  * @return #hipSuccess, #hipErrorInvalidResourceHandle
  */
 hipError_t hipStreamDestroy(hipStream_t stream) {
-    HIP_INIT_API(stream);
+    HIP_INIT_CB_API(hipStreamDestroy, stream);
 
     hipError_t e = hipSuccess;
 
@@ -214,7 +214,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
 
 //---
 hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
-    HIP_INIT_API(stream, flags);
+    HIP_INIT_CB_API(hipStreamGetFlags, stream, flags);
 
     if (flags == NULL) {
         return ihipLogStatus(hipErrorInvalidValue);
@@ -250,7 +250,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
 //---
 hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
                                 unsigned int flags) {
-    HIP_INIT_API(stream, callback, userData, flags);
+    HIP_INIT_CB_API(hipStreamAddCallback, stream, callback, userData, flags);
     hipError_t e = hipSuccess;
 
     // Create a thread in detached mode to handle callback