From ea7d5037b99c918851e26460bd33ca4c39c53503 Mon Sep 17 00:00:00 2001
From: Ioannis Assiouras <Ioannis.Assiouras@amd.com>
Date: Wed, 22 Feb 2023 00:05:49 +0000
Subject: [PATCH] SWDEV-381402 - Remove unused getNullStream() from device.
 Make stream destructor private.

Change-Id: Idde30a8bfe97a525bd9f9fb50698a5cb14b798fc


[ROCm/clr commit: 06927fd3c1f91680f92495abc4639678db9d043c]
---
 projects/clr/hipamd/src/hip_device.cpp         | 16 +---------------
 projects/clr/hipamd/src/hip_graph_internal.cpp |  2 +-
 projects/clr/hipamd/src/hip_internal.hpp       |  4 ++++
 projects/clr/hipamd/src/hip_mempool.cpp        |  6 +++---
 4 files changed, 9 insertions(+), 19 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp
index 092652d098..8782b6c35e 100644
--- a/projects/clr/hipamd/src/hip_device.cpp
+++ b/projects/clr/hipamd/src/hip_device.cpp
@@ -39,20 +39,6 @@ hip::Stream* Device::NullStream(bool skip_alloc) {
   return null_stream_;
 }
 
-// ================================================================================================
-hip::Stream* Device::GetNullStream() {
-  if (null_stream_ == nullptr) {
-    null_stream_ = new Stream(this, Stream::Priority::Normal, 0, true);
-  }
-
-  if (null_stream_ == nullptr) {
-    return nullptr;
-  }
-  // Wait for all active streams before executing commands on the default
-  iHipWaitActiveStreams(null_stream_);
-  return null_stream_;
-}
-
 // ================================================================================================
 bool Device::Create() {
   // Create default memory pool
@@ -152,7 +138,7 @@ Device::~Device() {
   }
 
   if (null_stream_!= nullptr) {
-    delete null_stream_;
+    null_stream_->release();
   }
 }
 
diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp
index 5733cc5c9d..f4060a0fe9 100644
--- a/projects/clr/hipamd/src/hip_graph_internal.cpp
+++ b/projects/clr/hipamd/src/hip_graph_internal.cpp
@@ -569,7 +569,7 @@ hipError_t hipGraphExec::Run(hipStream_t stream) {
       levelOrder_[0]->GetParentGraph()->FreeAllMemory();
     }
   }
-  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->GetNullStream()
+  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream()
                                         : reinterpret_cast<hip::Stream*>(stream);
   UpdateStream(parallelLists_, hip_stream, this);
   std::vector<amd::Command*> rootCommands;
diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp
index 0ebb4b5816..49a8832650 100644
--- a/projects/clr/hipamd/src/hip_internal.hpp
+++ b/projects/clr/hipamd/src/hip_internal.hpp
@@ -380,6 +380,10 @@ namespace hip {
       }
     }
     static bool existsActiveStreamForDevice(hip::Device* device);
+
+    /// The stream should be destroyed via release() rather than delete
+    private:
+      ~Stream() {};
   };
 
   /// HIP Device class
diff --git a/projects/clr/hipamd/src/hip_mempool.cpp b/projects/clr/hipamd/src/hip_mempool.cpp
index eea254f5a7..f798f8c813 100644
--- a/projects/clr/hipamd/src/hip_mempool.cpp
+++ b/projects/clr/hipamd/src/hip_mempool.cpp
@@ -70,7 +70,7 @@ hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) {
   if ((dev_ptr == nullptr) || (size == 0) || (!hip::isValid(stream))) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->GetNullStream() :
+  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() :
     reinterpret_cast<hip::Stream*>(stream);
   auto device = hip_stream->GetDevice();
   auto mem_pool = device->GetCurrentMemoryPool();
@@ -92,7 +92,7 @@ hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) {
   auto memory = getMemoryObject(dev_ptr, offset);
   if (memory != nullptr) {
     auto id = memory->getUserData().deviceId;
-    auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->GetNullStream() :
+    auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() :
       reinterpret_cast<hip::Stream*>(stream);
     if (!g_devices[id]->FreeMemory(memory, hip_stream)) {
       //! @todo It's not the most optimal logic. The current implementation has unconditional waits
@@ -241,7 +241,7 @@ hipError_t hipMallocFromPoolAsync(
   STREAM_CAPTURE(hipMallocAsync, stream, mem_pool, size, dev_ptr);
 
   auto mpool = reinterpret_cast<hip::MemoryPool*>(mem_pool);
-  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->GetNullStream() :
+  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() :
     reinterpret_cast<hip::Stream*>(stream);
   *dev_ptr = mpool->AllocateMemory(size, hip_stream);
   HIP_RETURN(hipSuccess);