From dc1f6f8b92315fb90d9694df85ae2ce7a4a4f7e0 Mon Sep 17 00:00:00 2001 From: Jon Chesterfield Date: Thu, 22 Jul 2021 10:29:30 +0100 Subject: [PATCH] [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 --- openmp/libomptarget/plugins/amdgpu/impl/internal.h | 16 +--------------- 1 file changed, 1 insertion(+), 15 deletions(-) diff --git a/openmp/libomptarget/plugins/amdgpu/impl/internal.h b/openmp/libomptarget/plugins/amdgpu/impl/internal.h index 21f0cc1..30dc393 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() { -- 2.7.4