Merge pull request #964 from eshcherb/roctracer-hip-frontend-190314

tracing callback layer update

hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -2924,7 +2924,7 @@ 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);
+const char* hipKernelNameRef(const hipFunction_t f);
 #ifdef __cplusplus
 } /* extern "C" */
 #endif

hipamd/src/hip_intercept.cpp

@@ -27,7 +27,7 @@ THE SOFTWARE.
 
 api_callbacks_table_t callbacks_table;
 
-extern std::string& FunctionSymbol(hipFunction_t f);
+extern std::string& FunctionSymbol(const 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) {

hipamd/src/hip_module.cpp

@@ -99,7 +99,7 @@ string ToString(hipFunction_t v) {
     return ss.str();
 };
 
-std::string& FunctionSymbol(hipFunction_t f) { return f->_name; };
+const std::string& FunctionSymbol(const hipFunction_t f) { return f->_name; };
 
 #define CHECK_HSA(hsaStatus, hipStatus)                                                            \
     if (hsaStatus != HSA_STATUS_SUCCESS) {                                                         \
@@ -262,7 +262,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     uint32_t localWorkSizeZ, size_t sharedMemBytes,
                                     hipStream_t hStream, void** kernelParams, void** extra,
                                     hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) {
-    HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
+    HIP_INIT_API(hipExtModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
                  localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra);
     return ihipLogStatus(ihipModuleLaunchKernel(
         f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,

hipamd/src/hip_peer.cpp

@@ -73,7 +73,7 @@ hipError_t ihipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_
  */
 //---
 hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
-    HIP_INIT_API(hipDeviceCanAccessPeer2, canAccessPeer, thisCtx, peerCtx);
+    HIP_INIT_API(NONE, canAccessPeer, thisCtx, peerCtx);
 
     return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, thisCtx, peerCtx));
 }
@@ -150,7 +150,7 @@ hipError_t ihipEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
 //---
 hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx,
                          size_t sizeBytes) {
-    HIP_INIT_API(hipMemcpyPeer2, dst, dstCtx, src, srcCtx, sizeBytes);
+    HIP_INIT_API(NONE, dst, dstCtx, src, srcCtx, sizeBytes);
 
     // TODO - move to ihip memory copy implementaion.
     // HCC has a unified memory architecture so device specifiers are not required.
@@ -161,7 +161,7 @@ hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t s
 //---
 hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice,
                               size_t sizeBytes, hipStream_t stream) {
-    HIP_INIT_API(hipMemcpyPeerAsync2, dst, dstDevice, src, srcDevice, sizeBytes, stream);
+    HIP_INIT_API(NONE, dst, dstDevice, src, srcDevice, sizeBytes, stream);
 
     // TODO - move to ihip memory copy implementaion.
     // HCC has a unified memory architecture so device specifiers are not required.