From 492ccc328158aa6ed9356969757bbebac0febca0 Mon Sep 17 00:00:00 2001
From: Sarbojit Sarkar <Sarbojit.Sarkar@amd.com>
Date: Thu, 1 Apr 2021 05:34:21 -0700
Subject: [PATCH] SWDEV-260454 - Fix for idenfying bad streams

Change-Id: I122ff6f47535c3c76ab56ba49ab1450ef886b15f
---
 include/hip/hip_runtime_api.h |  1 +
 rocclr/hip_internal.hpp       |  2 ++
 rocclr/hip_stream.cpp         | 44 ++++++++++++++++++++++++++++++++++-
 3 files changed, 46 insertions(+), 1 deletion(-)

diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h
index d8bb8c5b4e..6b9069df1d 100644
--- a/include/hip/hip_runtime_api.h
+++ b/include/hip/hip_runtime_api.h
@@ -267,6 +267,7 @@ typedef enum __HIP_NODISCARD hipError_t {
     hipErrorPeerAccessNotEnabled =
         705,  ///< Peer access was never enabled from the current device.
     hipErrorSetOnActiveProcess = 708,
+    hipErrorContextIsDestroyed = 709,
     hipErrorAssert = 710,  ///< Produced when the kernel calls assert.
     hipErrorHostMemoryAlreadyRegistered =
         712,  ///< Produced when trying to lock a page-locked memory.
diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp
index a22f848b9e..f5eac616e4 100755
--- a/rocclr/hip_internal.hpp
+++ b/rocclr/hip_internal.hpp
@@ -244,6 +244,8 @@ namespace hip {
   extern amd::HostQueue* getNullStream(amd::Context&);
   /// Get default stream of the thread
   extern amd::HostQueue* getNullStream();
+  /// Check if stream is valid
+  extern bool isValid(hipStream_t stream);
 };
 
 struct ihipExec_t {
diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp
index 322a5915a9..f4f6196a4d 100755
--- a/rocclr/hip_stream.cpp
+++ b/rocclr/hip_stream.cpp
@@ -123,7 +123,22 @@ void Stream::syncNonBlockingStreams() {
   }
 }
 
-};
+// ================================================================================================
+bool isValid(hipStream_t stream) {
+  // NULL stream is always valid
+  if (stream == nullptr) {
+    return true;
+  }
+
+  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+  amd::ScopedLock lock(streamSetLock);
+  if (streamSet.find(s) == streamSet.end()) {
+    return false;
+  }
+  return true;
+}
+
+};// hip namespace
 
 // ================================================================================================
 void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream) {
@@ -263,6 +278,9 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
   HIP_INIT_API(hipStreamGetFlags, stream, flags);
 
   if ((flags != nullptr) && (stream != nullptr)) {
+    if (!hip::isValid(stream)) {
+      return HIP_RETURN(hipErrorContextIsDestroyed);
+    }
     *flags = reinterpret_cast<hip::Stream*>(stream)->Flags();
   } else {
     HIP_RETURN(hipErrorInvalidValue);
@@ -275,6 +293,10 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
 hipError_t hipStreamSynchronize(hipStream_t stream) {
   HIP_INIT_API(hipStreamSynchronize, stream);
 
+  if (!hip::isValid(stream)) {
+    return HIP_RETURN(hipErrorContextIsDestroyed);
+  }
+
   // Wait for the current host queue
   hip::getQueue(stream)->finish();
 
@@ -289,6 +311,10 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
     HIP_RETURN(hipErrorInvalidHandle);
   }
 
+  if (!hip::isValid(stream)) {
+    return HIP_RETURN(hipErrorContextIsDestroyed);
+  }
+
   delete reinterpret_cast<hip::Stream*>(stream);
 
   HIP_RETURN(hipSuccess);
@@ -316,6 +342,10 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
     HIP_RETURN(hipErrorInvalidHandle);
   }
 
+  if (!hip::isValid(stream)) {
+    return HIP_RETURN(hipErrorContextIsDestroyed);
+  }
+
   amd::HostQueue* queue = hip::getQueue(stream);
 
   hip::Event* e = reinterpret_cast<hip::Event*>(event);
@@ -344,6 +374,10 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 hipError_t hipStreamQuery(hipStream_t stream) {
   HIP_INIT_API(hipStreamQuery, stream);
 
+  if (!hip::isValid(stream)) {
+    return HIP_RETURN(hipErrorContextIsDestroyed);
+  }
+
   amd::HostQueue* hostQueue = hip::getQueue(stream);
 
   amd::Command* command = hostQueue->getLastQueuedCommand(true);
@@ -369,6 +403,11 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
   if (callback == nullptr || flags != 0) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+
+  if (!hip::isValid(stream)) {
+    return HIP_RETURN(hipErrorContextIsDestroyed);
+  }
+
   amd::HostQueue* hostQueue = hip::getQueue(stream);
   amd::Command* last_command = hostQueue->getLastQueuedCommand(true);
   amd::Command::EventWaitList eventWaitList;
@@ -420,6 +459,9 @@ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize
 hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
   HIP_INIT_API(hipStreamGetPriority, stream, priority);
   if ((priority != nullptr) && (stream != nullptr)) {
+    if (!hip::isValid(stream)) {
+      return HIP_RETURN(hipErrorContextIsDestroyed);
+    }
     *priority = static_cast<int>(reinterpret_cast<hip::Stream*>(stream)->GetPriority());
   } else {
     HIP_RETURN(hipErrorInvalidValue);