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);