P4 to Git Change 1553709 by cpaquot@cpaquot-ocl-lc-lnx on 2018/05/11 11:24:30

SWDEV-145570 - [HIP] Make streamSet global and protect it
	By default from the spec, streamSet should be global and not per thread.
	There is a flag to make it per thread but we don't handle this yet. We
	would just add another variable that will be thread local and use it instead.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#12 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#10 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#11 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#8 edit


[ROCm/clr commit: c6b9701abb]

projects/clr/hipamd/api/hip/hip_context.cpp

@@ -64,12 +64,6 @@ void setCurrentContext(unsigned int index) {
   g_context = g_devices[index];
 }
 
-void syncStreams() {
-  for (const auto& it : streamSet) {
-    it->finish();
-  }
-}
-
 amd::HostQueue* getNullStream() {
   auto stream = g_nullStreams.find(getCurrentContext());
   if (stream == g_nullStreams.end()) {

projects/clr/hipamd/api/hip/hip_device_runtime.cpp

@@ -377,6 +377,8 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) {
 hipError_t hipDeviceSynchronize ( void ) {
   HIP_INIT_API();
 
+  hip::syncStreams();
+
   amd::HostQueue* queue = hip::getNullStream();
 
   if (!queue) {

projects/clr/hipamd/api/hip/hip_internal.hpp

@@ -61,7 +61,6 @@ namespace hip {
   extern void syncStreams();
 };
 extern std::vector<amd::Context*> g_devices;
-extern thread_local std::unordered_set<amd::HostQueue*> streamSet;
 extern hipError_t ihipDeviceGetCount(int* count);
 
 #endif // HIP_SRC_HIP_INTERNAL_H

projects/clr/hipamd/api/hip/hip_stream.cpp

@@ -22,8 +22,22 @@ THE SOFTWARE.
 
 #include <hip/hip_runtime.h>
 #include "hip_internal.hpp"
+#include "thread/monitor.hpp"
 
-thread_local std::unordered_set<amd::HostQueue*> streamSet;
+static amd::Monitor streamSetLock("Guards global stream set");
+static std::unordered_set<amd::HostQueue*> streamSet;
+
+namespace hip {
+
+void syncStreams() {
+  amd::ScopedLock lock(streamSetLock);
+
+  for (const auto& it : streamSet) {
+    it->finish();
+  }
+}
+
+};
 
 static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   amd::Device* device = hip::getCurrentContext()->devices()[0];
@@ -38,7 +52,11 @@ static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int fl
 
   if (!(flags & hipStreamNonBlocking)) {
     hip::syncStreams();
-    streamSet.insert(queue);
+
+    {
+      amd::ScopedLock lock(streamSetLock);
+      streamSet.insert(queue);
+    }
   }
 
   *stream = reinterpret_cast<hipStream_t>(as_cl(queue));
@@ -100,6 +118,8 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
     return hipErrorInvalidResourceHandle;
   }
 
+  amd::ScopedLock lock(streamSetLock);
+
   amd::HostQueue* hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   hostQueue->release();
   streamSet.erase(hostQueue);