diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp
index a93f49f5f7..29d83e43fb 100644
--- a/hipamd/src/hip_hcc.cpp
+++ b/hipamd/src/hip_hcc.cpp
@@ -87,7 +87,6 @@ unsigned g_numLogicalThreads;
 std::atomic<int> g_lastShortTid(1);
 
 
-thread_local ShortTid tls_shortTid;
 
 /*
  Implementation of malloc and free device functions.
@@ -175,10 +174,12 @@ __device__ void* __hip_hc_free(void *ptr)
 
 // This is the implicit context used by all HIP commands.
 // It can be set by hipSetDevice or by the CTX manipulation commands:
-
 thread_local hipError_t tls_lastHipError = hipSuccess;
 
 
+thread_local ShortTid tls_shortTid;
+
+
 
 
 //=================================================================================================
@@ -238,7 +239,8 @@ hipError_t ihipSynchronize(void)
 //=================================================================================================
 // ihipStream_t:
 //=================================================================================================
-ShortTid::ShortTid() 
+ShortTid::ShortTid()  :
+    _apiSeqNum(0)
 { 
     _shortTid = g_lastShortTid.fetch_add(1); 
 
@@ -273,8 +275,6 @@ ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int
     };
 
 
-
-
     tprintf(DB_SYNC, " streamCreate: stream=%p\n", this);
 };
 
@@ -1292,7 +1292,8 @@ void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, c
 {
     if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {
         std::stringstream os;
-        os  << "<<hip-api: hipLaunchKernel '" << kernelName << "'"
+        os  << "<<hip-api tid:" << tls_shortTid.tid() << "." << tls_shortTid.incApiSeqNum()
+            << " hipLaunchKernel '" << kernelName << "'"
             << " gridDim:"  << lp->grid_dim
             << " groupDim:" << lp->group_dim
             << " sharedMem:+" << lp->dynamic_group_mem_bytes
diff --git a/hipamd/src/hip_hcc.h b/hipamd/src/hip_hcc.h
index 19338642fe..668011a8c8 100644
--- a/hipamd/src/hip_hcc.h
+++ b/hipamd/src/hip_hcc.h
@@ -66,10 +66,15 @@ public:
 
     ShortTid() ;
 
-    int tid() { return _shortTid; };
-private:
+    int      tid()       const { return _shortTid; };
+    uint64_t incApiSeqNum() { return ++_apiSeqNum; };
+    uint64_t apiSeqNum() const { return _apiSeqNum; };
 
-    int _shortTid;
+private:
+    int      _shortTid;
+
+    // monotonically increasing API sequence number for this threa.
+    uint64_t _apiSeqNum; 
 };
 
 //---
@@ -145,7 +150,7 @@ extern const char *API_COLOR_END;
     if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {\
         std::string s = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\
         if (COMPILE_HIP_DB && HIP_TRACE_API) {\
-            fprintf (stderr, "%s<<hip-api:tid:%d %s\n%s" , API_COLOR, tls_shortTid.tid(), s.c_str(), API_COLOR_END);\
+            fprintf (stderr, "%s<<hip-api tid:%d.%lu %s\n%s" , API_COLOR, tls_shortTid.tid(), tls_shortTid.incApiSeqNum(), s.c_str(), API_COLOR_END);\
         }\
         SCOPED_MARKER(s.c_str(), "HIP", NULL);\
     }\
@@ -175,7 +180,7 @@ extern const char *API_COLOR_END;
         tls_lastHipError = localHipStatus;\
         \
         if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API) {\
-            fprintf(stderr, "  %ship-api:tid:%d %-30s ret=%2d (%s)>>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, tls_shortTid.tid(), __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\
+            fprintf(stderr, "  %ship-api tid:%d.%lu %-30s ret=%2d (%s)>>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, tls_shortTid.tid(),tls_shortTid.apiSeqNum(),  __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\
         }\
         localHipStatus;\
     })