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

[JonChesterfield]

Summary

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.

[jdoerfert]

LG. If we don't need the pool, maybe remove it?

[JonChesterfield]

It's a performance thing. I'm not sure signal creation is expensive enough to be worth caching, particularly without a thread local cache, but haven't checked.

They might be cheap to create on gfx9 and expensive on gfx8 or similar. Will need to trace through HSA to see how many operations are avoided by the cache.

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

 namespace core {
 class TaskgroupImpl;
 class TaskImpl;
 class Kernel;
 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() {
     size_t N = state.size();
     for (size_t i = 0; i < N; i++) {
       hsa_signal_t signal = state.front();
       state.pop();
       hsa_status_t rc = hsa_signal_destroy(signal);