From ddbd7039b033164053f4aba504297736cb4bfd9d Mon Sep 17 00:00:00 2001 From: Rahul Manocha Date: Mon, 19 Aug 2024 14:20:46 -0700 Subject: [PATCH] SWDEV-478921 - Destroy Queue created by Coop Launch Change-Id: I7f31ce05421479ff1de138cae26aafa071e956e2 --- rocclr/device/rocm/rocdevice.cpp | 8 +++++++- rocclr/device/rocm/rocdevice.hpp | 2 +- rocclr/device/rocm/rocvirtual.cpp | 2 +- 3 files changed, 9 insertions(+), 3 deletions(-) diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index bfd46a2491..78c55eeb9b 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -3221,7 +3221,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, return queue; } -void Device::releaseQueue(hsa_queue_t* queue, const std::vector& cuMask) { +void Device::releaseQueue(hsa_queue_t* queue, const std::vector& cuMask, bool coop_queue) { for (auto& it : cuMask.size() == 0 ? queuePool_ : queueWithCUMaskPool_) { auto qIter = it.find(queue); if (qIter != it.end()) { @@ -3248,6 +3248,12 @@ void Device::releaseQueue(hsa_queue_t* queue, const std::vector& cuMas } } } + if (coop_queue) { // cooperative queue + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting CG enabled hardware queue %p ", + queue->base_address); + hsa_queue_destroy(queue); + } + } void* Device::getOrCreateHostcallBuffer(hsa_queue_t* queue, bool coop_queue, diff --git a/rocclr/device/rocm/rocdevice.hpp b/rocclr/device/rocm/rocdevice.hpp index 2f91df4a16..dc968b2ee4 100644 --- a/rocclr/device/rocm/rocdevice.hpp +++ b/rocclr/device/rocm/rocdevice.hpp @@ -563,7 +563,7 @@ class Device : public NullDevice { amd::CommandQueue::Priority priority = amd::CommandQueue::Priority::Normal); //! Release HSA queue - void releaseQueue(hsa_queue_t*, const std::vector& cuMask = {}); + void releaseQueue(hsa_queue_t*, const std::vector& cuMask = {}, bool coop_queue = false); //! For the given HSA queue, return an existing hostcall buffer or create a //! new one. queuePool_ keeps a mapping from HSA queue to hostcall buffer. diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 632a058f2b..cbd4c99b00 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -1328,7 +1328,7 @@ VirtualGPU::~VirtualGPU() { } if (gpu_queue_) { - roc_device_.releaseQueue(gpu_queue_, cuMask_); + roc_device_.releaseQueue(gpu_queue_, cuMask_, cooperative_); } }