diff --git a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h
--- a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h
+++ b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.h
@@ -142,6 +142,8 @@
                               void *data, uint32_t private_segment_size,
                               uint32_t group_segment_size, hsa_queue_t **queue);
 
+hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);
+
 uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);
 
 uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
diff --git a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp
--- a/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/dynamic_hsa/hsa.cpp
@@ -29,6 +29,7 @@
 DLWRAP(hsa_signal_store_screlease, 2);
 DLWRAP(hsa_signal_wait_scacquire, 5);
 DLWRAP(hsa_queue_create, 8);
+DLWRAP(hsa_queue_destroy, 1);
 DLWRAP(hsa_queue_load_read_index_scacquire, 1);
 DLWRAP(hsa_queue_add_write_index_relaxed, 2);
 DLWRAP(hsa_memory_copy, 3);
diff --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -797,6 +797,7 @@
           UINT32_MAX, UINT32_MAX, &HSAQueues[i]);
       if (rc != HSA_STATUS_SUCCESS) {
         DP("Failed to create HSA queue %d\n", i);
+        HSAQueues[i] = nullptr;
         return;
       }