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/impl/internal.h b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
--- a/openmp/libomptarget/plugins/amdgpu/impl/internal.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
@@ -80,17 +80,23 @@
   SignalPoolT() {}
   SignalPoolT(const SignalPoolT &) = delete;
   SignalPoolT(SignalPoolT &&) = delete;
-  ~SignalPoolT() {
+  ~SignalPoolT() {}
+
+  hsa_status_t clear() {
+    // Drop contents before shutting down hsa
     size_t N = state.size();
+    hsa_status_t res = HSA_STATUS_SUCCESS;
     for (size_t i = 0; i < N; i++) {
       hsa_signal_t signal = state.front();
       state.pop();
       hsa_status_t rc = hsa_signal_destroy(signal);
       if (rc != HSA_STATUS_SUCCESS) {
-        DP("Signal pool destruction failed\n");
+        res = rc;
       }
     }
+    return res;
   }
+
   size_t size() {
     lock l(&mutex);
     return state.size();
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
@@ -772,6 +772,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;
       }
 
@@ -810,7 +811,7 @@
     hostrpc_terminate();
 
     hsa_status_t Err;
-    for (uint32_t I = 0; I < HSAExecutables.size(); I++) {
+    for (size_t I = 0; I < HSAExecutables.size(); I++) {
       Err = hsa_executable_destroy(HSAExecutables[I]);
       if (Err != HSA_STATUS_SUCCESS) {
         DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
@@ -818,6 +819,18 @@
       }
     }
 
+    Err = FreeSignalPool.clear();
+    if (Err != HSA_STATUS_SUCCESS) {
+      DP("Error destroying signal pool: %s\n", Err);
+    }
+
+    for (size_t I = 0; I < HSAQueues.size(); I++) {
+      Err = hsa_queue_destroy(HSAQueues[I]);
+      if (Err != HSA_STATUS_SUCCESS) {
+        DP("Error destroying hsa queue: %s\n", get_error_string(Err));
+      }
+    }
+
     Err = hsa_shut_down();
     if (Err != HSA_STATUS_SUCCESS) {
       DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Shutting down HSA",