diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h
index cb2612ffbe..9553a170c9 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h
@@ -330,7 +330,14 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {
 
   hsa_queue_t* public_handle_;
 
+  /// Next available queue id.
+  uint64_t GetQueueId() { return hsa_queue_counter_++; }
+
  private:
+
+  // HSA Queue ID - used to bind a unique ID
+  static std::atomic<uint64_t> hsa_queue_counter_;
+
   DISALLOW_COPY_AND_ASSIGN(Queue);
 };
 }   //  namespace core
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h
index 3e1112de13..35fa50dfa6 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h
@@ -251,11 +251,6 @@ class Runtime {
   /// set.
   hsa_status_t GetSystemInfo(hsa_system_info_t attribute, void* value);
 
-  /// @brief Query next available queue id.
-  ///
-  /// @retval Next available queue id.
-  uint32_t GetQueueId();
-
   /// @brief Register a callback function @p handler that is associated with
   /// @p signal to asynchronous event monitor thread.
   ///
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
index ca5e3deb64..3a4352fc70 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
@@ -225,9 +225,12 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
   // Complete populating the doorbell signal structure.
   signal_.legacy_hardware_doorbell_ptr =
       (volatile uint32_t*)queue_rsrc.Queue_DoorBell;
-  // Complete populating the amd_queue_ structure.
-  amd_queue_.hsa_queue.id = queue_id_ = queue_rsrc.QueueId;
 
+  // Bind Id of Queue such that is unique i.e. it is not re-used by another
+  // queue (AQL, HOST) in the same process during its lifetime.
+  amd_queue_.hsa_queue.id = this->GetQueueId();
+
+  queue_id_ = queue_rsrc.QueueId;
   MAKE_NAMED_SCOPE_GUARD(QueueGuard, [&]() { hsaKmtDestroyQueue(queue_id_); });
 
   MAKE_NAMED_SCOPE_GUARD(EventGuard, [&]() {
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/host_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/host_queue.cpp
index 2eab97b3b2..aaee7933ee 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/host_queue.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/host_queue.cpp
@@ -77,7 +77,7 @@ HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t
   amd_queue_.hsa_queue.base_address = ring_;
   amd_queue_.hsa_queue.size = size_;
   amd_queue_.hsa_queue.doorbell_signal = doorbell_signal;
-  amd_queue_.hsa_queue.id = queue_count_++;
+  amd_queue_.hsa_queue.id = this->GetQueueId();
   amd_queue_.hsa_queue.type = type;
   amd_queue_.hsa_queue.features = features;
 #ifdef HSA_LARGE_MODEL
diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp
index 6c8a661839..a86bbf8158 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp
@@ -46,6 +46,9 @@
 namespace rocr {
 namespace core {
 
+// HSA Queue ID - used to bind a unique ID
+std::atomic<uint64_t> Queue::hsa_queue_counter_(0);
+
 void Queue::DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data) {
   if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) {
     const char* msg = "UNKNOWN ERROR";