[Openmp-commits] [openmp] dc1f6f8 - [libomptarget][amdgpu][nfc] Drop dead signal pool setup

[Jon Chesterfield via Openmp-commits]

Author: Jon Chesterfield
Date: 2021-07-22T10:29:32+01:00
New Revision: dc1f6f8b92315fb90d9694df85ae2ce7a4a4f7e0

URL: https://github.com/llvm/llvm-project/commit/dc1f6f8b92315fb90d9694df85ae2ce7a4a4f7e0
DIFF: https://github.com/llvm/llvm-project/commit/dc1f6f8b92315fb90d9694df85ae2ce7a4a4f7e0.diff

LOG: [libomptarget][amdgpu][nfc] Drop dead signal pool setup

This class is instantiated once in rtl.cpp before hsa_init is
called. The hsa_signal_create call therefore fails leaving the pool empty.

This signal pool is a legacy from ATMI where it was constructed after hsa_init.
Moving the state into the rtl.cpp global class disabled the initial populating
of the pool without noticeably changing performance. Just rechecked with a fix
that allocates the signals after hsa_init and that also doesn't noticeably
change performance.

This patch therefore drops the initialisation. Only change from main is to
drop a DEBUG_PRINT statement that would say the pool initial size is zero.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D106515

Added: 
    

Modified: 
    openmp/libomptarget/plugins/amdgpu/impl/internal.h

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/plugins/amdgpu/impl/internal.h b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
index 21f0cc15ad7cb..30dc393870b55 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/internal.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
@@ -99,21 +99,7 @@ class KernelImpl;
 } // namespace core
 
 struct SignalPoolT {
-  SignalPoolT() {
-    // If no signals are created, and none can be created later,
-    // will ultimately fail at pop()
-
-    unsigned N = 1024; // default max pool size from atmi
-    for (unsigned i = 0; i < N; i++) {
-      hsa_signal_t new_signal;
-      hsa_status_t err = hsa_signal_create(0, 0, NULL, &new_signal);
-      if (err != HSA_STATUS_SUCCESS) {
-        break;
-      }
-      state.push(new_signal);
-    }
-    DEBUG_PRINT("Signal Pool Initial Size: %lu\n", state.size());
-  }
+  SignalPoolT() {}
   SignalPoolT(const SignalPoolT &) = delete;
   SignalPoolT(SignalPoolT &&) = delete;
   ~SignalPoolT() {