https://gcc.gnu.org/g:45d083d8c21d286d1da0202ec152f8130d5c8dc2
commit 45d083d8c21d286d1da0202ec152f8130d5c8dc2 Author: Andrew Stubbs <[email protected]> Date: Tue Feb 17 17:47:01 2026 +0000 libgomp, amdgcn: Fix issues with dynamic OpenMP thread scaling This patch, based on one originally applied to the OG11 branch in 2021, fixes a bug in which testcases using thread_limit larger than the number of physical threads would crash with a memory fault. This was exacerbated in testcases with a lot of register pressure because the autoscaling reduces the number of physical threads to compensate for the increased resource usage. We specifically saw this happen in the t-reduction testcase in the external omptests testsuite. With this patch that testcase now passes, and a couple of other failures are also fixed. The included test case was greatly reduced from the t-reduction testcase with c-vise and hand-editing. The code is nonsensical, but it was triggering the memory fault with only 13 threads. It was also checked with nvidia offloading and on x86_64 without offloading. libgomp/ChangeLog: * config/gcn/bar.h (gomp_barrier_init): Limit thread count to the actual physical number. Return count, not void. (gomp_barrier_reinit): Likewise. * config/gcn/team.c (gomp_team_start): Don't attempt to set up threads that do not exist. * config/linux/bar.h (gomp_barrier_init): Return count, not void. (gomp_barrier_reinit): Likewise. * config/nvptx/bar.h (gomp_barrier_init): Likewise. (gomp_barrier_reinit): Likewise. * config/posix/bar.c (gomp_barrier_init): Likewise. (gomp_barrier_reinit): Likewise. * config/posix/bar.h: (gomp_barrier_init): Likewise. (gomp_barrier_reinit): Likewise. * config/posix/simple-bar.h: (gomp_barrier_init): Likewise. (gomp_barrier_reinit): Likewise. * config/rtems/bar.h: (gomp_barrier_init): Likewise. (gomp_barrier_reinit): Likewise. * team.c (gomp_new_team): Set nthreads to actual rather than requested number of threads. * testsuite/libgomp.c/thread-limit-6.c: New test case. Co-Authored-by: Sandra Loosemore <[email protected]> Diff: --- libgomp/config/gcn/bar.h | 14 ++++++-- libgomp/config/gcn/team.c | 4 +++ libgomp/config/linux/bar.h | 8 +++-- libgomp/config/nvptx/bar.h | 8 +++-- libgomp/config/posix/bar.c | 6 ++-- libgomp/config/posix/bar.h | 4 +-- libgomp/config/posix/simple-bar.h | 8 ++--- libgomp/config/rtems/bar.h | 8 +++-- libgomp/team.c | 2 +- libgomp/testsuite/libgomp.c/thread-limit-6.c | 52 ++++++++++++++++++++++++++++ 10 files changed, 97 insertions(+), 17 deletions(-) diff --git a/libgomp/config/gcn/bar.h b/libgomp/config/gcn/bar.h index 6e838ff54a89..11c0404b3090 100644 --- a/libgomp/config/gcn/bar.h +++ b/libgomp/config/gcn/bar.h @@ -53,18 +53,28 @@ typedef unsigned int gomp_barrier_state_t; #define BAR_CANCELLED 4 #define BAR_INCR 8 -static inline void gomp_barrier_init (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_init (gomp_barrier_t *bar, unsigned count) { + unsigned actual_thread_count = __builtin_gcn_dim_size (1); + if (count > actual_thread_count) + count = actual_thread_count; bar->total = count; bar->awaited = count; bar->awaited_final = count; bar->generation = 0; + return count; } -static inline void gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) { + unsigned actual_thread_count = __builtin_gcn_dim_size (1); + if (count > actual_thread_count) + count = actual_thread_count; __atomic_add_fetch (&bar->awaited, count - bar->total, MEMMODEL_RELAXED); bar->total = count; + return count; } static inline void gomp_barrier_destroy (gomp_barrier_t *bar) diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c index be148c3f8bf6..fdf3f3d63958 100644 --- a/libgomp/config/gcn/team.c +++ b/libgomp/config/gcn/team.c @@ -243,6 +243,10 @@ gomp_team_start (void (*fn) (void *), void *data, unsigned nthreads, .taskgroup = taskgroup }; + unsigned actual_thread_count = __builtin_gcn_dim_size (1); + if (nthreads > actual_thread_count) + nthreads = actual_thread_count; + if (nthreads != 1) { /* When there's more than one thread, we expect that we're operating on diff --git a/libgomp/config/linux/bar.h b/libgomp/config/linux/bar.h index 4dc0d3cca994..9e999a1dba11 100644 --- a/libgomp/config/linux/bar.h +++ b/libgomp/config/linux/bar.h @@ -53,18 +53,22 @@ typedef unsigned int gomp_barrier_state_t; #define BAR_CANCELLED 4 #define BAR_INCR 8 -static inline void gomp_barrier_init (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_init (gomp_barrier_t *bar, unsigned count) { bar->total = count; bar->awaited = count; bar->awaited_final = count; bar->generation = 0; + return count; } -static inline void gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) { __atomic_add_fetch (&bar->awaited, count - bar->total, MEMMODEL_ACQ_REL); bar->total = count; + return count; } static inline void gomp_barrier_destroy (gomp_barrier_t *bar) diff --git a/libgomp/config/nvptx/bar.h b/libgomp/config/nvptx/bar.h index aa2592ba5b3f..010b7b146659 100644 --- a/libgomp/config/nvptx/bar.h +++ b/libgomp/config/nvptx/bar.h @@ -51,18 +51,22 @@ typedef unsigned int gomp_barrier_state_t; #define BAR_CANCELLED 4 #define BAR_INCR 8 -static inline void gomp_barrier_init (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_init (gomp_barrier_t *bar, unsigned count) { bar->total = count; bar->awaited = count; bar->awaited_final = count; bar->generation = 0; + return count; } -static inline void gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) { __atomic_add_fetch (&bar->awaited, count - bar->total, MEMMODEL_ACQ_REL); bar->total = count; + return count; } static inline void gomp_barrier_destroy (gomp_barrier_t *bar) diff --git a/libgomp/config/posix/bar.c b/libgomp/config/posix/bar.c index c46659bd2645..202519c5d337 100644 --- a/libgomp/config/posix/bar.c +++ b/libgomp/config/posix/bar.c @@ -31,7 +31,7 @@ #include "libgomp.h" -void +unsigned int gomp_barrier_init (gomp_barrier_t *bar, unsigned count) { gomp_mutex_init (&bar->mutex1); @@ -44,6 +44,7 @@ gomp_barrier_init (gomp_barrier_t *bar, unsigned count) bar->arrived = 0; bar->generation = 0; bar->cancellable = false; + return count; } void @@ -61,12 +62,13 @@ gomp_barrier_destroy (gomp_barrier_t *bar) gomp_sem_destroy (&bar->sem2); } -void +unsigned int gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) { gomp_mutex_lock (&bar->mutex1); bar->total = count; gomp_mutex_unlock (&bar->mutex1); + return count; } void diff --git a/libgomp/config/posix/bar.h b/libgomp/config/posix/bar.h index 026daca793d5..8fd673cb42ac 100644 --- a/libgomp/config/posix/bar.h +++ b/libgomp/config/posix/bar.h @@ -58,8 +58,8 @@ typedef unsigned int gomp_barrier_state_t; #define BAR_CANCELLED 4 #define BAR_INCR 8 -extern void gomp_barrier_init (gomp_barrier_t *, unsigned); -extern void gomp_barrier_reinit (gomp_barrier_t *, unsigned); +extern unsigned int gomp_barrier_init (gomp_barrier_t *, unsigned); +extern unsigned int gomp_barrier_reinit (gomp_barrier_t *, unsigned); extern void gomp_barrier_destroy (gomp_barrier_t *); extern void gomp_barrier_wait (gomp_barrier_t *); diff --git a/libgomp/config/posix/simple-bar.h b/libgomp/config/posix/simple-bar.h index 9452a2d5fdc8..e2e19fc17b55 100644 --- a/libgomp/config/posix/simple-bar.h +++ b/libgomp/config/posix/simple-bar.h @@ -36,16 +36,16 @@ typedef struct gomp_barrier_t bar; } gomp_simple_barrier_t; -static inline void +static inline unsigned int gomp_simple_barrier_init (gomp_simple_barrier_t *bar, unsigned count) { - gomp_barrier_init (&bar->bar, count); + return gomp_barrier_init (&bar->bar, count); } -static inline void +static inline unsigned int gomp_simple_barrier_reinit (gomp_simple_barrier_t *bar, unsigned count) { - gomp_barrier_reinit (&bar->bar, count); + return gomp_barrier_reinit (&bar->bar, count); } static inline void diff --git a/libgomp/config/rtems/bar.h b/libgomp/config/rtems/bar.h index 80fb1cd3be87..3fd81d31e150 100644 --- a/libgomp/config/rtems/bar.h +++ b/libgomp/config/rtems/bar.h @@ -54,19 +54,23 @@ typedef unsigned int gomp_barrier_state_t; #define BAR_CANCELLED 4 #define BAR_INCR 8 -static inline void gomp_barrier_init (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_init (gomp_barrier_t *bar, unsigned count) { bar->total = count; bar->awaited = count; bar->awaited_final = count; bar->generation = 0; _Futex_Initialize (&bar->futex); + return count; } -static inline void gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) +static inline unsigned int +gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count) { __atomic_add_fetch (&bar->awaited, count - bar->total, MEMMODEL_ACQ_REL); bar->total = count; + return count; } static inline void gomp_barrier_destroy (gomp_barrier_t *bar) diff --git a/libgomp/team.c b/libgomp/team.c index 5282a3133ba9..f1496533ddc9 100644 --- a/libgomp/team.c +++ b/libgomp/team.c @@ -187,7 +187,7 @@ gomp_new_team (unsigned nthreads) #ifndef HAVE_SYNC_BUILTINS gomp_mutex_init (&team->work_share_list_free_lock); #endif - gomp_barrier_init (&team->barrier, nthreads); + nthreads = gomp_barrier_init (&team->barrier, nthreads); gomp_mutex_init (&team->task_lock); team->nthreads = nthreads; diff --git a/libgomp/testsuite/libgomp.c/thread-limit-6.c b/libgomp/testsuite/libgomp.c/thread-limit-6.c new file mode 100644 index 000000000000..ffe7395808c6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/thread-limit-6.c @@ -0,0 +1,52 @@ +/* { dg-additional-options "-O3" } */ + +/* This test is reduced from a larger example that used to give a + "Memory access fault" on AMD GCN due to creating too many threads. */ + +int main() { + char Ac_0; + short As[2871]; + int Ai[2871]; + int Ci_0; + int Bi_0; + int Ai_0; + long All[2871]; + long Cll_0; + long Bll_0; + long All_0; + float Af_0; + double Cd_0; + double Bd_0; + double Ad_0; + short Rs7, Rs10; + int Ri1, Ri2, Ri5, Ri6, Ri7, Ri10; + long Rll1, Rll2, Rll5, Rll6, Rll7, Rll10; + float Rf1; + double Rd2; + + int i = 0; + for (int threads = 0; threads < 512; threads++) + { +#pragma omp target teams num_teams(1) thread_limit(1024) +#pragma omp parallel if (threads > 1) num_threads(threads) + { + Rs7 |= 1 < 0; + Rs10 = Rs10 || As[i] > 0; + Ri1 += Ai_0 + 0; + Ri2 += Ai_0 + (Bi_0 + Ci_0); + Ri5 *= i % 1000 == 0 ? 2 : 1; + Ri6 &= 0; + Ri7 |= 1 < 0; + Ri10 = Ri10 || Ai[i] > 0; + Rll1 += 0; + Rll2 += All_0 + (Bll_0 + Cll_0); + Rll5 *= i % 1000 == 0 ? : 1; + Rll6 &= 0; + Rll7 |= 1ll < 0; + Rll10 = Rll10 || All[i] > 0; + Rf1 += Af_0 + 0; + Rd2 += Ad_0 + (Bd_0 + Cd_0); + } + } + return 0; +}
