diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md
index f61c3b106a..d54032c6df 100644
--- a/projects/clr/hipamd/README.md
+++ b/projects/clr/hipamd/README.md
@@ -32,6 +32,7 @@ HIP releases are typically of two types. The tag naming convention is different
 - [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
 - [HIP Porting Guide](docs/markdown/hip_porting_guide.md)
 - [HIP Porting Driver Guide](docs/markdown/hip_porting_driver_api.md)
+- [HIP Profiling and Debugging](docs/markdown/hip_profiling.md)
 - [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
 - [hipify-clang](hipify-clang/README.md)
 - [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
diff --git a/projects/clr/hipamd/docs/markdown/hip_profiling.md b/projects/clr/hipamd/docs/markdown/hip_profiling.md
index 6e5cde700d..b5c4672464 100644
--- a/projects/clr/hipamd/docs/markdown/hip_profiling.md
+++ b/projects/clr/hipamd/docs/markdown/hip_profiling.md
@@ -267,6 +267,11 @@ info: check result
 PASSED!
 ```
 
+HIP_TRACE_API supports multiple levels of debug information:
+   - 0x1 = print all HIP APIs
+   - 0x2 = print HIP APIs which initiate GPU kernels, copies, or memsets.  Includes hipLaunchKernel, hipMemcpy*, hipMemset*.
+   - 0x4 = print HIP APIs which allocate or free memory.  Includes hipMalloc, hipHostMalloc, hipFree, hipHostFree.
+
 
 #### Color
 Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.
diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp
index a655e35aa1..07604fe85d 100644
--- a/projects/clr/hipamd/src/hip_hcc.cpp
+++ b/projects/clr/hipamd/src/hip_hcc.cpp
@@ -1435,7 +1435,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
 void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
 {
 
-    if ((HIP_TRACE_API & (1<<TRACE_CMD)) || HIP_PROFILE_API || (COMPILE_HIP_DB && HIP_TRACE_API)) {
+    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:";
diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h
index 132f099ce8..7787242ca7 100644
--- a/projects/clr/hipamd/src/hip_hcc_internal.h
+++ b/projects/clr/hipamd/src/hip_hcc_internal.h
@@ -188,9 +188,10 @@ extern const char *API_COLOR_END;
 
 //---
 //HIP Trace modes
-#define TRACE_ALL 0 // 0x1
-#define TRACE_CMD 1 // 0x2
-#define TRACE_MEM 2 // 0x4
+#define TRACE_ALL  0 // 0x1
+#define TRACE_KCMD 1 // 0x2, kernel command
+#define TRACE_MCMD 2 // 0x4, memory command
+#define TRACE_MEM  3 // 0x8
 
 
 //---
@@ -275,12 +276,13 @@ extern void recordApiTrace(std::string *fullStr, const std::string &apiStr);
     API_TRACE(0, __VA_ARGS__);
 
 
-// Like above, but will trace with DB_CMD.
-// Replace HIP_INIT_API with this call inside important APIs that launch work on the GPU:
+// Like above, but will trace with TRACE_CMD.  
+// Replace HIP_INIT_API with this call inside HIP APIs that launch work on the GPU:
 // kernel launches, copy commands, memory sets, etc.
-#define HIP_INIT_CMD_API(...) \
+#define HIP_INIT_SPECIAL_API(tbit, ...) \
     HIP_INIT()\
-    API_TRACE((HIP_TRACE_API&(1<<TRACE_CMD)), __VA_ARGS__);
+    API_TRACE((HIP_TRACE_API&(1<<tbit)), __VA_ARGS__);
+
 
 // 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 hipGetLastError,
diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index c4f0f64e50..cef676b572 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -213,7 +213,7 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi
 
 hipError_t hipMalloc(void** ptr, size_t sizeBytes)
 {
-    HIP_INIT_API(ptr, sizeBytes);
+    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, sizeBytes);
     HIP_SET_DEVICE();
     hipError_t  hip_status = hipSuccess;
 
@@ -244,7 +244,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
 
 hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 {
-    HIP_INIT_CMD_API(ptr, sizeBytes, flags);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), ptr, sizeBytes, flags);
     HIP_SET_DEVICE();
     hipError_t hip_status = hipSuccess;
 
@@ -308,7 +308,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags)
 // width in bytes
 hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height)
 {
-    HIP_INIT_CMD_API(ptr, pitch, width, height);
+    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, pitch, width, height);
     HIP_SET_DEVICE();
     hipError_t  hip_status = hipSuccess;
 
@@ -349,7 +349,7 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel
 hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
         size_t width, size_t height, unsigned int flags)
 {
-    HIP_INIT_CMD_API(array, desc, width, height, flags);
+    HIP_INIT_SPECIAL_API((TRACE_MEM), array, desc, width, height, flags);
     HIP_SET_DEVICE();
     hipError_t  hip_status = hipSuccess;
 
@@ -490,7 +490,7 @@ hipError_t hipHostUnregister(void *hostPtr)
 
 hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
 {
-    HIP_INIT_CMD_API(symbolName, src, count, offset, kind);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind);
 
     if(symbolName == nullptr)
     {
@@ -560,7 +560,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_CMD_API(symbolName, src, count, offset, kind, stream);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind, stream);
 
     if(symbolName == nullptr)
     {
@@ -637,7 +637,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_CMD_API(dst, src, sizeBytes, kind);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -659,7 +659,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_CMD_API(dst, src, sizeBytes);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -681,7 +681,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_CMD_API(dst, src, sizeBytes);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -703,7 +703,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_CMD_API(dst, src, sizeBytes);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -725,7 +725,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte
 
 hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes)
 {
-    HIP_INIT_CMD_API(dst, src, sizeBytes);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -750,7 +750,7 @@ 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_CMD_API(dst, src, sizeBytes, kind, stream);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind, stream);
 
     return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream));
 
@@ -759,21 +759,21 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
 
 hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream)
 {
-    HIP_INIT_CMD_API(dst, src, sizeBytes, stream);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
 
     return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyHostToDevice, stream));
 }
 
 hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
 {
-    HIP_INIT_CMD_API(dst, src, sizeBytes, stream);
+    HIP_INIT_SPECIAL_API((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_CMD_API(dst, src, sizeBytes, stream);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
 
     return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream));
 }
@@ -782,7 +782,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
         size_t width, size_t height, hipMemcpyKind kind) {
 
-    HIP_INIT_CMD_API(dst, dpitch, src, spitch, width, height, kind);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
 
     if(width > dpitch || width > spitch)
         return ihipLogStatus(hipErrorUnknown);
@@ -826,7 +826,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
 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_CMD_API(dst, wOffset, hOffset, src, spitch, width, height, kind);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -879,7 +879,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_CMD_API(dst, wOffset, hOffset, src, count, kind);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, count, kind);
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
@@ -938,7 +938,7 @@ ihipMemsetKernel(hipStream_t stream,
 // 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_CMD_API(dst, value, sizeBytes, stream);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes, stream);
 
     hipError_t e = hipSuccess;
 
@@ -988,7 +988,7 @@ hipError_t hipMemsetAsync(void* dst, int  value, size_t sizeBytes, hipStream_t s
 
 hipError_t hipMemset(void* dst, int  value, size_t sizeBytes )
 {
-    HIP_INIT_CMD_API(dst, value, sizeBytes);
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
 
     hipError_t e = hipSuccess;
 
@@ -1148,7 +1148,7 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size)
 
 hipError_t hipFree(void* ptr)
 {
-    HIP_INIT_API(ptr);
+    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
 
     hipError_t hipStatus = hipErrorInvalidDevicePointer;
 
@@ -1176,7 +1176,7 @@ hipError_t hipFree(void* ptr)
 
 hipError_t hipHostFree(void* ptr)
 {
-    HIP_INIT_API(ptr);
+    HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
 
     // Synchronize to ensure all work has finished.
     ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
@@ -1210,7 +1210,7 @@ hipError_t hipFreeHost(void* ptr)
 
 hipError_t hipFreeArray(hipArray* array)
 {
-    HIP_INIT_API(array);
+    HIP_INIT_SPECIAL_API((TRACE_MEM), array);
 
     hipError_t hipStatus = hipErrorInvalidDevicePointer;