From ecde1f5b960a31ce9e7156bb587c7d1bbe574d6a Mon Sep 17 00:00:00 2001
From: Sarbojit Sarkar <Sarbojit.Sarkar@amd.com>
Date: Tue, 21 Sep 2021 09:25:16 +0000
Subject: [PATCH] SWDEV-292547 - Fix for hipStreamPerThread

Change-Id: I9328c7ec92d954433f8e69732d91e65fa6fe013a


[ROCm/clr commit: 9e4766d99008b5a20ace4323218ecaacf7c5e122]
---
 projects/clr/hipamd/src/hip_stream.cpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp
index 583770ec8a..22b509c439 100644
--- a/projects/clr/hipamd/src/hip_stream.cpp
+++ b/projects/clr/hipamd/src/hip_stream.cpp
@@ -292,11 +292,14 @@ public:
   ~stream_per_thread() {
     if (m_stream != nullptr && hip::isValid(m_stream)) {
       delete reinterpret_cast<hip::Stream*>(m_stream);
+      m_stream = nullptr;
     }
   }
 
   hipStream_t& get() {
-    if (m_stream == nullptr) {
+    // There is a scenario where hipResetDevice destroys stream per thread
+    // hence isValid check is required to make sure only valid stream is used
+    if (m_stream == nullptr || !hip::isValid(m_stream)) {
       ihipStreamCreate(&m_stream, hipStreamDefault, hip::Stream::Priority::Normal);
     }
     return m_stream;