SWDEV-478921 - Destroy Queue created by Coop Launch

Change-Id: I7f31ce05421479ff1de138cae26aafa071e956e2

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<uint32_t>& cuMask) {
+void Device::releaseQueue(hsa_queue_t* queue, const std::vector<uint32_t>& 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<uint32_t>& 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,

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<uint32_t>& cuMask = {});
+  void releaseQueue(hsa_queue_t*, const std::vector<uint32_t>& 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.

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