diff --git a/projects/hip/rocclr/hip_internal.hpp b/projects/hip/rocclr/hip_internal.hpp index 248b127d4b..a22f848b9e 100755 --- a/projects/hip/rocclr/hip_internal.hpp +++ b/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& 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; diff --git a/projects/hip/rocclr/hip_stream.cpp b/projects/hip/rocclr/hip_stream.cpp index 339b61acb2..49475bca7d 100755 --- a/projects/hip/rocclr/hip_stream.cpp +++ b/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(stream)->Destroy(); + delete reinterpret_cast(stream); HIP_RETURN(hipSuccess); }