consolidate thread local storage (#915)

* all thread local access now through single struct

* clean up old commented-out code, more use of GET_TLS()

* fewer calls to GET_TLS by passing tls as a funtion argument

* revert unnecessary change to printf

* fix failing tests due to TLS change

* fix merge conflicts in ihipOccupancyMaxActiveBlocksPerMultiprocessor

hipamd/src/hip_hcc.cpp

@@ -139,24 +139,12 @@ std::atomic<int> g_lastShortTid(1);
 std::vector<ProfTrigger> g_dbStartTriggers;
 std::vector<ProfTrigger> g_dbStopTriggers;
 
-//=================================================================================================
-// Thread-local storage:
-//=================================================================================================
-
-// 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 TidInfo tls_tidInfo;
-
-
 //=================================================================================================
 // Top-level "free" functions:
 //=================================================================================================
-uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr) {
-    auto apiSeqNum = tls_tidInfo.apiSeqNum();
-    auto tid = tls_tidInfo.tid();
+uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::string& apiStr) {
+    auto apiSeqNum = tls->tidInfo.apiSeqNum();
+    auto tid = tls->tidInfo.tid();
 
     if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) {
         printf("info: resume profiling at %lu\n", apiSeqNum);
@@ -179,7 +167,7 @@ uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr) {
 
 
     if (COMPILE_HIP_DB && HIP_TRACE_API) {
-        fprintf(stderr, "%s<<hip-api pid:%d tid:%s @%lu%s\n", API_COLOR, tls_tidInfo.pid(), fullStr->c_str(), apiStartTick,
+        fprintf(stderr, "%s<<hip-api pid:%d tid:%s @%lu%s\n", API_COLOR, tls->tidInfo.pid(), fullStr->c_str(), apiStartTick,
                 API_COLOR_END);
     }
 
@@ -206,32 +194,18 @@ ihipCtx_t* ihipGetPrimaryCtx(unsigned deviceIndex) {
     return device ? device->getPrimaryCtx() : NULL;
 };
 
-
-static thread_local ihipCtx_t* tls_defaultCtx = nullptr;
-void ihipSetTlsDefaultCtx(ihipCtx_t* ctx) { tls_defaultCtx = ctx; }
-
-
-//---
-// TODO - review the context creation strategy here.  Really should be:
-//  - first "non-device" runtime call creates the context for this thread.  Allowed to call
-//  setDevice first.
-//  - hipDeviceReset destroys the primary context for device?
-//  - Then context is created again for next usage.
-ihipCtx_t* ihipGetTlsDefaultCtx() {
-    // Per-thread initialization of the TLS:
-    if ((tls_defaultCtx == nullptr) && (g_deviceCnt > 0)) {
-        ihipSetTlsDefaultCtx(ihipGetPrimaryCtx(0));
-    }
-    return tls_defaultCtx;
-}
-
-hipError_t ihipSynchronize(void) {
+hipError_t ihipSynchronize(TlsData *tls) {
     ihipGetTlsDefaultCtx()->locked_waitAllStreams();  // ignores non-blocking streams, this waits
                                                       // for all activity to finish.
 
     return (hipSuccess);
 }
 
+TlsData* tls_get_ptr() {
+    static thread_local TlsData data;
+    return &data;
+}
+
 //=================================================================================================
 // ihipStream_t:
 //=================================================================================================
@@ -1482,7 +1456,7 @@ hipError_t hip_init() {
 }
 }
 
-hipError_t ihipStreamSynchronize(hipStream_t stream) {
+hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) {
     hipError_t e = hipSuccess;
 
     if (stream == hipStreamNull) {
@@ -1503,7 +1477,8 @@ void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) {
     // Synchronize stream
     tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n",
             ToString(cb->_stream).c_str());
-    e = ihipStreamSynchronize(cb->_stream);
+    GET_TLS();
+    e = ihipStreamSynchronize(tls, cb->_stream);
 
     // Call registered callback function
     cb->_callback(cb->_stream, e, cb->_userData);
@@ -1518,6 +1493,7 @@ void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) {
 hipStream_t ihipSyncAndResolveStream(hipStream_t stream, bool lockAcquired) {
     if (stream == hipStreamNull) {
         // Submitting to NULL stream, call locked_syncDefaultStream to wait for all other streams:
+        GET_TLS();
         ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
         tprintf(DB_SYNC, "ihipSyncAndResolveStream %s wait on default stream\n",
                 ToString(stream).c_str());
@@ -1581,15 +1557,16 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp,
                            const hipStream_t stream) {
     if ((HIP_TRACE_API & (1 << TRACE_KCMD)) || HIP_PROFILE_API ||
         (COMPILE_HIP_DB & HIP_TRACE_API)) {
+        GET_TLS();
         std::stringstream os;
-        os << tls_tidInfo.pid() << " " << tls_tidInfo.tid() << "." << tls_tidInfo.apiSeqNum() << " hipLaunchKernel '"
+        os << tls->tidInfo.pid() << " " << tls->tidInfo.tid() << "." << tls->tidInfo.apiSeqNum() << " hipLaunchKernel '"
            << kernelName << "'"
            << " gridDim:" << lp->grid_dim << " groupDim:" << lp->group_dim << " sharedMem:+"
            << lp->dynamic_group_mem_bytes << " " << *stream;
 
         if (COMPILE_HIP_DB && HIP_TRACE_API) {
             std::string fullStr;
-            recordApiTrace(&fullStr, os.str());
+            recordApiTrace(tls, &fullStr, os.str());
         }
 
         if (HIP_PROFILE_API == 0x1) {