SWDEV-381402 - Remove unused getNullStream() from device. Make stream destructor private.

Change-Id: Idde30a8bfe97a525bd9f9fb50698a5cb14b798fc


[ROCm/clr commit: 06927fd3c1]

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

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;

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

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