AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add per-thread API seqnum to debug

Browse Source 
Change-Id: Ib13733a3e84cd56bae13a32bae40f936c20b7543
This commit is contained in:
Ben Sander
2016-10-24 20:24:12 -05:00
parent f8b356dce8
commit 45e2dd4d69
2 changed files with 17 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_hcc.cpp
+7 -6
View File
@@ -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
hipamd/src/hip_hcc.h
+10 -5
View File
@@ -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;\
})