SWDEV-277543 - fix destruction of the stream object

Change-Id: If2f532e66099805a7f05181e63f5454893a1a288


[ROCm/hip commit: 7600326eb2]

projects/hip/rocclr/hip_internal.hpp

@@ -147,6 +147,7 @@ namespace hip {
   public:
     Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false,
            const std::vector<uint32_t>& cuMask = {});
+    ~Stream();
 
     /// Creates the hip stream object, including AMD host queue
     bool Create();
@@ -154,7 +155,6 @@ namespace hip {
     /// Get device AMD host queue object. The method can allocate the queue
     amd::HostQueue* asHostQueue(bool skip_alloc = false);
 
-    void Destroy();
     void Finish() const;
     /// Get device ID associated with the current stream;
     int DeviceId() const;

projects/hip/rocclr/hip_stream.cpp

@@ -36,6 +36,17 @@ Stream::Stream(hip::Device* dev, Priority p,
   : queue_(nullptr), lock_("Stream Callback lock"), device_(dev),
     priority_(p), flags_(f), null_(null_stream), cuMask_(cuMask) {}
 
+// ================================================================================================
+Stream::~Stream() {
+  if (queue_ != nullptr) {
+    amd::ScopedLock lock(streamSetLock);
+    streamSet.erase(this);
+
+    queue_->release();
+    queue_ = nullptr;
+  }
+}
+
 // ================================================================================================
 bool Stream::Create() {
   // Enable queue profiling if a profiler is attached which sets the callback_table flag
@@ -68,10 +79,10 @@ bool Stream::Create() {
     amd::ScopedLock lock(streamSetLock);
     streamSet.insert(this);
     queue_ = queue;
-  } else {
-    queue_ = queue;
-    Destroy();
+  } else if (queue != nullptr) {
+    queue->release();
   }
+
   return result;
 }
 
@@ -91,18 +102,6 @@ amd::HostQueue* Stream::asHostQueue(bool skip_alloc) {
   return queue_;
 }
 
-// ================================================================================================
-void Stream::Destroy() {
-  if (queue_ != nullptr) {
-    amd::ScopedLock lock(streamSetLock);
-    streamSet.erase(this);
-
-    queue_->release();
-    queue_ = nullptr;
-  }
-  delete this;
-}
-
 // ================================================================================================
 void Stream::Finish() const {
   if (queue_ != nullptr) {
@@ -192,6 +191,7 @@ static hipError_t ihipStreamCreate(hipStream_t* stream,
   hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags, false, cuMask);
 
   if (hStream == nullptr || !hStream->Create()) {
+    delete hStream;
     return hipErrorOutOfMemory;
   }
 
@@ -286,7 +286,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
     HIP_RETURN(hipErrorInvalidHandle);
   }
 
-  reinterpret_cast<hip::Stream*>(stream)->Destroy();
+  delete reinterpret_cast<hip::Stream*>(stream);
 
   HIP_RETURN(hipSuccess);
 }