Make ihipDevice_t thread-safe.

Move critical data into separate class and protect with LockAccessor
wrapper class.

For device, the streams list is the critical data since it is modified when
streams are created or destroyed.   The streams list is accessed in
several places including when synchronizing across all streams on the
device (ie from the default stream).
Other device data is set once by the device cosntructor and is not critical
so

All functions which acquire the LockAccessor now named with "locked_" prefix.


[ROCm/hip commit: 530ab9434a]

projects/hip/src/hip_device.cpp

@@ -150,8 +150,7 @@ hipError_t hipDeviceSynchronize(void)
 {
     HIP_INIT_API();
 
-    ihipGetTlsDefaultDevice()->waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
-
+    ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
 
     return ihipLogStatus(hipSuccess);
 }
@@ -174,15 +173,12 @@ hipError_t hipDeviceReset(void)
 
     if (device) {
         //---
-        //Wait for pending activity to complete?
-        //TODO - check if this is required behavior:
-        for (auto streamI=device->_streams.begin(); streamI!=device->_streams.end(); streamI++) {
-            ihipStream_t *stream = *streamI;
-            stream->wait();
-        }
+        //Wait for pending activity to complete?  TODO - check if this is required behavior:
+        
+        device->locked_waitAllStreams();
 
         // Release device resources (streams and memory):
-        device->reset(); 
+        device->locked_reset(); 
     }
 
     return ihipLogStatus(hipSuccess);

projects/hip/src/hip_event.cpp

@@ -67,7 +67,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
             // TODO-HCC fix this - is CUDA this conservative or still uses device timestamps?
             // TODO-HCC can we use barrier or event marker to implement better solution?
             ihipDevice_t *device = ihipGetTlsDefaultDevice();
-            device->syncDefaultStream(true);
+            device->locked_syncDefaultStream(true);
 
             eh->_timestamp = hc::get_system_ticks();
             eh->_state = hipEventStatusRecorded;
@@ -117,7 +117,7 @@ hipError_t hipEventSynchronize(hipEvent_t event)
             return ihipLogStatus(hipSuccess);
         } else if (eh->_stream == NULL) {
             ihipDevice_t *device = ihipGetTlsDefaultDevice();
-            device->syncDefaultStream(true);
+            device->locked_syncDefaultStream(true);
             return ihipLogStatus(hipSuccess);
         } else {
 #if __hcc_workweek__ >= 16033

projects/hip/src/hip_hcc.cpp

@@ -145,6 +145,8 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, SeqNu
 ihipStream_t::~ihipStream_t()
 {
     _signalPool.clear();
+    // Hack to catch memory issues, in particular accesses to _acc after it has been destroyed.
+    //memset (&_av, 0x0, sizeof(hc::accelerator_view&));
 }
 
 
@@ -348,6 +350,8 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna
             hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (_last_kernel_future.get_native_handle()));
             if (hsaSignal) {
                 *waitSignal = * hsaSignal;
+            } else {
+                assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail.  Confirm this never happens.
             }
         } else if (_last_copy_signal) {
             needSync = 1;
@@ -383,10 +387,13 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna
 //
 //Reset the device - this is called from hipDeviceReset.
 //Device may be reset multiple times, and may be reset after init.
-void ihipDevice_t::reset()
+void ihipDevice_t::locked_reset()
 {
+    // Obtain mutex access to the device critical data, release by destructor
+    Locked_ihipDeviceCritical_t  l(_criticalData);
+
     // Reset and remove streams:
-    _streams.clear();
+    l->streams().clear();
 
     // Reset and release all memory stored in the tracker:
     am_memtracker_reset(_acc);
@@ -395,7 +402,7 @@ void ihipDevice_t::reset()
 
 
 //---
-void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned flags)
+void ihipDevice_t::locked_init(unsigned device_index, hc::accelerator acc, unsigned flags)
 {
     _stream_id = 0;
 
@@ -416,8 +423,12 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned fla
 
     getProperties(&_props);
 
-    _default_stream = new ihipStream_t(device_index, acc.get_default_view(), _stream_id++, hipStreamDefault);
-    this->_streams.push_back(_default_stream);
+    { 
+        Locked_ihipDeviceCritical_t  l(_criticalData);
+        _default_stream = new ihipStream_t(device_index, acc.get_default_view(), _stream_id++, hipStreamDefault);
+        l->streams().push_back(_default_stream);
+    }
+
     tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream);
 
 
@@ -670,11 +681,13 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop)
 // Implement "default" stream syncronization
 //   This waits for all other streams to drain before continuing.
 //   If waitOnSelf is set, this additionally waits for the default stream to empty.
-void ihipDevice_t::syncDefaultStream(bool waitOnSelf)
+void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf)
 {
+    Locked_ihipDeviceCritical_t  l(_criticalData);
+
     tprintf(DB_SYNC, "syncDefaultStream\n");
 
-    for (auto streamI=_streams.begin(); streamI!=_streams.end(); streamI++) {
+    for (auto streamI=l->const_streams().begin(); streamI!=l->const_streams().end(); streamI++) {
         ihipStream_t *stream = *streamI;
 
         // Don't wait for streams that have "opted-out" of syncing with NULL stream.
@@ -690,13 +703,31 @@ void ihipDevice_t::syncDefaultStream(bool waitOnSelf)
     }
 }
 
+//---
+void ihipDevice_t::locked_addStream(ihipStream_t *s)
+{
+    Locked_ihipDeviceCritical_t  l(_criticalData);
+
+    l->streams().push_back(s);
+}
+
+//---
+void ihipDevice_t::locked_removeStream(ihipStream_t *s)
+{
+    Locked_ihipDeviceCritical_t  l(_criticalData);
+
+    l->streams().remove(s);
+}
+
 
 //---
 //Heavyweight synchronization that waits on all streams, ignoring hipStreamNonBlocking flag.
-void ihipDevice_t::waitAllStreams()
+void ihipDevice_t::locked_waitAllStreams()
 {
+    Locked_ihipDeviceCritical_t  l(_criticalData);
+
     tprintf(DB_SYNC, "waitAllStream\n");
-    for (auto streamI=_streams.begin(); streamI!=_streams.end(); streamI++) {
+    for (auto streamI=l->const_streams().begin(); streamI!=l->const_streams().end(); streamI++) {
         (*streamI)->wait();
     }
 }
@@ -864,7 +895,7 @@ void ihipInit()
                 //If device is not in visible devices list, ignore
                 continue;
             }
-            g_devices[g_deviceCnt].init(g_deviceCnt, accs[i], hipDeviceMapHost);
+            g_devices[g_deviceCnt].locked_init(g_deviceCnt, accs[i], hipDeviceMapHost);
             g_deviceCnt++;
         }
     }
@@ -933,7 +964,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
         ihipDevice_t *device = ihipGetTlsDefaultDevice();
 
 #ifndef HIP_API_PER_THREAD_DEFAULT_STREAM
-        device->syncDefaultStream(false);
+        device->locked_syncDefaultStream(false);
 #endif
         return device->_default_stream;
     } else {
@@ -1302,6 +1333,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a
 // TODO - describe naming convention. ihip _.  No accessors.  No early returns from functions. Set status to success at top, only set error codes in implementation.  No tabs.
 //        Caps convention _ or camelCase
 //        if { }
+//        Should use ihip* data structures inside code rather than app-facing hip.  For example, use ihipDevice_t (rather than hipDevice_t), ihipStream_t (rather than hipStream_t).
 // TODO - describe MT strategy
 //
 //// TODO - add identifier numbers for streams and devices to help with debugging.

projects/hip/src/hip_memory.cpp

@@ -459,7 +459,7 @@ hipError_t hipFree(void* ptr)
     hipError_t hipStatus = hipErrorInvalidDevicePointer;
 
    // Synchronize to ensure all work has finished.
-    ihipGetTlsDefaultDevice()->waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
+    ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
 
     if (ptr) {
         hc::accelerator acc;

projects/hip/src/hip_stream.cpp

@@ -42,7 +42,9 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
     //This matches CUDA stream behavior:
 
     auto istream = new ihipStream_t(device->_device_index, acc.create_view(), device->_stream_id, flags);
-    device->_streams.push_back(istream);
+
+    device->locked_addStream(istream);
+
     *stream = istream;
     tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream);
 
@@ -82,7 +84,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
 
     if (stream == NULL) {
         ihipDevice_t *device = ihipGetTlsDefaultDevice();
-        device->syncDefaultStream(true/*waitOnSelf*/);
+        device->locked_syncDefaultStream(true/*waitOnSelf*/);
     } else {
         stream->wait();
         e = hipSuccess;
@@ -106,7 +108,7 @@ hipError_t hipStreamDestroy(hipStream_t stream)
     //--- Drain the stream:
     if (stream == NULL) {
         ihipDevice_t *device = ihipGetTlsDefaultDevice();
-        device->syncDefaultStream(true/*waitOnSelf*/);
+        device->locked_syncDefaultStream(true/*waitOnSelf*/);
     } else {
         stream->wait();
         e = hipSuccess;
@@ -115,7 +117,7 @@ hipError_t hipStreamDestroy(hipStream_t stream)
     ihipDevice_t *device = stream->getDevice();
 
     if (device) {
-        device->_streams.remove(stream);
+        device->locked_removeStream(stream);
         delete stream;
     } else {
         e = hipErrorInvalidResourceHandle;