Merge 'master' into 'amd-master'

Change-Id: I510b01084b5aad8382044fe828a099896fa12aff

src/grid_launch.cpp

@@ -72,8 +72,8 @@ namespace hip_impl
             }
 
             if (COMPILE_HIP_DB && HIP_TRACE_API) {
-                std::cerr << API_COLOR << os.str() << API_COLOR_END
-                          << std::endl;
+                std::string fullStr;
+                recordApiTrace(&fullStr, os.str());
             }
         }
     }

src/hip_device.cpp

@@ -358,7 +358,8 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device
 
 hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
 {
-    HIP_INIT_API(name,len, device);
+    // Cast to void* here to avoid printing garbage in debug modes.
+    HIP_INIT_API((void*)name,len, device);
     hipError_t e = hipSuccess;
     auto deviceHandle = ihipGetDevice(device);
     int nameLen = strlen(deviceHandle->_props.name);
@@ -369,7 +370,8 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
 
 hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
 {
-    HIP_INIT_API(pciBusId, len, device);
+    // Cast to void* here to avoid printing garbage in debug modes.
+    HIP_INIT_API((void*)pciBusId, len, device);
     hipError_t e = hipErrorInvalidValue;
     int deviceCount = 0;
     ihipGetDeviceCount( &deviceCount );

src/hip_hcc.cpp

@@ -1429,9 +1429,7 @@ void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, c
 {
 
     if ((HIP_TRACE_API & (1<<TRACE_KCMD)) || HIP_PROFILE_API || (COMPILE_HIP_DB & HIP_TRACE_API)) {
-        std::stringstream os_pre;
         std::stringstream os;
-        os_pre  << "<<hip-api tid:";
         os  << tls_tidInfo.tid() << "." << tls_tidInfo.apiSeqNum()
             << " hipLaunchKernel '" << kernelName << "'"
             << " gridDim:"  << lp->grid_dim
@@ -1439,6 +1437,11 @@ void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, c
             << " sharedMem:+" << lp->dynamic_group_mem_bytes
             << " " << *stream;
 
+        if (COMPILE_HIP_DB && HIP_TRACE_API) {
+            std::string fullStr;
+            recordApiTrace(&fullStr, os.str());
+        }
+
         if (HIP_PROFILE_API == 0x1) {
             std::string shortAtpString("hipLaunchKernel:");
             shortAtpString += kernelName;
@@ -1446,10 +1449,6 @@ void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, c
         } else if (HIP_PROFILE_API == 0x2) {
             MARKER_BEGIN(os.str().c_str(), "HIP");
         }
-
-        if (COMPILE_HIP_DB && HIP_TRACE_API) {
-            std::cerr << API_COLOR << os.str() << API_COLOR_END << std::endl;
-        }
     }
 }
 

src/hip_memory.cpp

@@ -879,7 +879,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
     return ihipLogStatus(e);
 }
 
-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);
     if(width > dpitch || width > spitch)