From c6b9701abb0d7bccc507b03f2012b82c62a97343 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Fri, 11 May 2018 11:34:01 -0400
Subject: [PATCH] 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
---
 hipamd/api/hip/hip_context.cpp        |  6 ------
 hipamd/api/hip/hip_device_runtime.cpp |  2 ++
 hipamd/api/hip/hip_internal.hpp       |  1 -
 hipamd/api/hip/hip_stream.cpp         | 24 ++++++++++++++++++++++--
 4 files changed, 24 insertions(+), 9 deletions(-)

diff --git a/hipamd/api/hip/hip_context.cpp b/hipamd/api/hip/hip_context.cpp
index 8559b641fc..e75c8e5b5e 100644
--- a/hipamd/api/hip/hip_context.cpp
+++ b/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()) {
diff --git a/hipamd/api/hip/hip_device_runtime.cpp b/hipamd/api/hip/hip_device_runtime.cpp
index c84abe02da..5b60623761 100644
--- a/hipamd/api/hip/hip_device_runtime.cpp
+++ b/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) {
diff --git a/hipamd/api/hip/hip_internal.hpp b/hipamd/api/hip/hip_internal.hpp
index 1f97be451a..197e607597 100644
--- a/hipamd/api/hip/hip_internal.hpp
+++ b/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
diff --git a/hipamd/api/hip/hip_stream.cpp b/hipamd/api/hip/hip_stream.cpp
index 7cb359b1b9..f4995bcd81 100644
--- a/hipamd/api/hip/hip_stream.cpp
+++ b/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);