[gcc/devel/omp/gcc-16] libgomp, amdgcn: Fix issues with dynamic OpenMP thread scaling

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;
+}