openmp: Re-add support for private references to variable-length arrays with allocato...

The previous patch for PR113436 fixed the testsuite regressions, but disabled
support for allocators when applied to references to variable-length objects
in private clauses.  This patch re-adds it.

2026-02-28  Kwok Cheung Yeung  <kcyeung@baylibre.com>

gcc/

	PR middle-end/113436
	* omp-low.cc (lower_omp_target): Merge branches for allocating memory
	for private clauses.  Add handling for references when allocator
	clause not specified.

gcc/testsuite/

	PR middle-end/113436
	* g++.dg/gomp/pr113436.C: Rename to...
	* g++.dg/gomp/pr113436-1.C: ... this.  Remove restriction on C++
	dialect.
	(f): Remove use of auto.
	* g++.dg/gomp/pr113436-2.C: New.  Original renamed to...
	* g++.dg/gomp/pr113436-5.C: ... this.  Add tests for alignment.
	(f): Test references to VLAs of pointers.
	* g++.dg/gomp/pr113436-3.C: New.
	* g++.dg/gomp/pr113436-4.C: New.

libgomp/

	PR middle-end/113436
	* testsuite/libgomp.c++/pr113436-1.C (test_vla_by_ref): New.
	(main): Add call to test_vla_by_ref.
	* testsuite/libgomp.c++/pr113436-2.C (test_vla_by_ref): New.
	(main): Add call to test_vla_by_ref.

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index aeed1d25e8af4bffe866a320ebb6486925b8ade1..b93012107f19d8942516074a813a582994d2bb4f 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -14329,7 +14329,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
        switch (OMP_CLAUSE_CODE (c))
          {
-           tree var;
+           tree var, new_var, *allocate_ptr;
          default:
            break;
          case OMP_CLAUSE_MAP:
@@ -14458,10 +14458,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
          case OMP_CLAUSE_PRIVATE:
            var = OMP_CLAUSE_DECL (c);
            by_ref = omp_privatize_by_reference (var);
-           if (is_variable_sized (var))
+           new_var = lookup_decl (var, ctx);
+           allocate_ptr = alloc_map.get (new_var);
+           if (is_variable_sized (var, by_ref)
+               || (!allocate_ptr && by_ref && !is_gimple_omp_oacc (ctx->stmt)))
              {
-               tree new_var = lookup_decl (var, ctx);
-               tree *allocate_ptr = alloc_map.get (new_var);
                if (allocate_ptr)
                  {
                    gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
@@ -14482,8 +14483,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                if (!allocate_ptr)
                  {
                    tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
-                   tree al = size_int (DECL_ALIGN (var));
-                   x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
+                   tree ty = TREE_TYPE (new_var);
+                   if (by_ref)
+                     ty = TREE_TYPE (ty);
+                   x = TYPE_SIZE_UNIT (ty);
+                   if (TREE_CONSTANT (x))
+                     break;
+                   tree al = size_int (TYPE_ALIGN (ty));
                    x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
                    x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
                  }
@@ -14493,27 +14499,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                gimple_seq_add_stmt (&new_body,
                                     gimple_build_assign (new_pvar, x));
              }
-           else if (by_ref && !is_gimple_omp_oacc (ctx->stmt))
-             {
-               location_t clause_loc = OMP_CLAUSE_LOCATION (c);
-               tree new_var = lookup_decl (var, ctx);
-               tree x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
-               if (TREE_CONSTANT (x))
-                 break;
-               else
-                 {
-                   tree atmp
-                     = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
-                   tree rtype = TREE_TYPE (TREE_TYPE (new_var));
-                   tree al = size_int (TYPE_ALIGN (rtype));
-                   x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
-                 }
-
-               x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
-               gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
-               gimple_seq_add_stmt (&new_body,
-                                    gimple_build_assign (new_var, x));
-             }
            break;
          case OMP_CLAUSE_FIRSTPRIVATE:
            var = OMP_CLAUSE_DECL (c);

diff --git a/gcc/testsuite/g++.dg/gomp/pr113436.C b/gcc/testsuite/g++.dg/gomp/pr113436-1.C
similarity index 94%
rename from gcc/testsuite/g++.dg/gomp/pr113436.C
rename to gcc/testsuite/g++.dg/gomp/pr113436-1.C
index 7078d88ccffd4b879502ad8f84616193eec00a0d..60920c6ab61db0d2f96f3c1e3e6b7f4d9a236236 100644 (file)
--- a/gcc/testsuite/g++.dg/gomp/pr113436.C
+++ b/gcc/testsuite/g++.dg/gomp/pr113436-1.C
@@ -1,6 +1,6 @@
 // PR middle-end/113436
 // { dg-do "compile" }
-// { dg-options "-std=gnu++20 -fopenmp -fdump-tree-omplower" }
+// { dg-options "-fopenmp -fdump-tree-omplower" }
 
 // #include <omp.h>
 typedef __UINTPTR_TYPE__ omp_uintptr_t;
@@ -31,7 +31,7 @@ typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM
 void f()
 {
   int a[10];
-  auto &aRef = a;
+  int (&aRef)[10] = a;
 
   #pragma omp target firstprivate(aRef) \
                     allocate(align(128), allocator(omp_low_lat_mem_alloc): aRef)

diff --git a/gcc/testsuite/g++.dg/gomp/pr113436-2.C b/gcc/testsuite/g++.dg/gomp/pr113436-2.C
index 70f6ffc196e69a736d13dfb48fae9c12b6fc0167..a681a2cce745438d133ea1c348cc959a5da86dde 100644 (file)
--- a/gcc/testsuite/g++.dg/gomp/pr113436-2.C
+++ b/gcc/testsuite/g++.dg/gomp/pr113436-2.C
@@ -2,16 +2,43 @@
 // { dg-do "compile" }
 // { dg-options "-fopenmp -fdump-tree-omplower" }
 
+// #include <omp.h>
+typedef __UINTPTR_TYPE__ omp_uintptr_t;
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM
+{
+  omp_null_allocator = 0,
+  omp_default_mem_alloc = 1,
+  omp_large_cap_mem_alloc = 2,
+  omp_const_mem_alloc = 3,
+  omp_high_bw_mem_alloc = 4,
+  omp_low_lat_mem_alloc = 5,
+  omp_cgroup_mem_alloc = 6,
+  omp_pteam_mem_alloc = 7,
+  omp_thread_mem_alloc = 8,
+  ompx_gnu_pinned_mem_alloc = 200,
+  ompx_gnu_managed_mem_alloc = 201,
+  __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+
 void f(int x)
 {
   int a[x];
-  int (&c)[x] = a;
+  int (&aRef)[x] = a;
 
-  #pragma omp target private (c)
-  {
-    c[0] = 1;
-  }
+  #pragma omp target private(aRef) \
+                    allocate(align(128), allocator(omp_high_bw_mem_alloc): aRef)
+    aRef[0] = 1;
 }
 
-// Ensure that the size of memory allocated for the VLA is from a variable rather than a constant.
-// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, \[0-9\]\+\\\);" "omplower" } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128, D\\\.\[0-9\]\+, 4\\\);" "omplower" } }
+// { dg-final { scan-tree-dump "aRef = D\\\.\[0-9\]\+;" "omplower" } }
+// { dg-final { scan-tree-dump "\\\(\\\*aRef\\\)\\\[0\\\] = 1;" "omplower" } }
+// { dg-final { scan-tree-dump "__builtin_GOMP_free \\\(D\\\.\[0-9\]\+, 4\\\);" "omplower" } }

diff --git a/gcc/testsuite/g++.dg/gomp/pr113436-3.C b/gcc/testsuite/g++.dg/gomp/pr113436-3.C
new file mode 100644 (file)
index 0000000..c2dbdc8
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/pr113436-3.C
@@ -0,0 +1,55 @@
+// PR middle-end/113436
+// { dg-do "compile" }
+// { dg-options "-fopenmp -fdump-tree-omplower" }
+
+// #include <omp.h>
+typedef __UINTPTR_TYPE__ omp_uintptr_t;
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM
+{
+  omp_null_allocator = 0,
+  omp_default_mem_alloc = 1,
+  omp_large_cap_mem_alloc = 2,
+  omp_const_mem_alloc = 3,
+  omp_high_bw_mem_alloc = 4,
+  omp_low_lat_mem_alloc = 5,
+  omp_cgroup_mem_alloc = 6,
+  omp_pteam_mem_alloc = 7,
+  omp_thread_mem_alloc = 8,
+  ompx_gnu_pinned_mem_alloc = 200,
+  ompx_gnu_managed_mem_alloc = 201,
+  __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+void f(int x)
+{
+  int a[x], *b[x];
+  int (&aRef)[x] = a;
+  int *(&bRef)[x] = b;
+
+  #pragma omp target firstprivate (aRef, bRef) \
+                    allocate (allocator (omp_low_lat_mem_alloc): aRef) \
+                    allocate (allocator (omp_large_cap_mem_alloc): bRef)
+  {
+    aRef[0] = 0;
+    bRef[0] = 0;
+  }
+}
+
+// Ensure that the size of memory allocated for the VLA is from a variable
+// rather than a constant, and that the default alignments are as expected.
+
+// { dg-final { scan-tree-dump "a\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "b\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target ilp32 } } }
+// { dg-final { scan-tree-dump "b\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 64\\\);" "omplower" { target { lp64 || llp64 } } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(4, D\\\.\[0-9\]\+, 5\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(4, D\\\.\[0-9\]\+, 2\\\);" "omplower" { target ilp32 } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(8, D\\\.\[0-9\]\+, 2\\\);" "omplower" { target { lp64 || llp64 } } } }
+// { dg-final { scan-tree-dump "aRef = D\\\.\[0-9\]\+;" "omplower" } }
+// { dg-final { scan-tree-dump "bRef = D\\\.\[0-9\]\+;" "omplower" } }

diff --git a/gcc/testsuite/g++.dg/gomp/pr113436-4.C b/gcc/testsuite/g++.dg/gomp/pr113436-4.C
new file mode 100644 (file)
index 0000000..03e2568
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/pr113436-4.C
@@ -0,0 +1,55 @@
+// PR middle-end/113436
+// { dg-do "compile" }
+// { dg-options "-fopenmp -fdump-tree-omplower" }
+
+// #include <omp.h>
+typedef __UINTPTR_TYPE__ omp_uintptr_t;
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM
+{
+  omp_null_allocator = 0,
+  omp_default_mem_alloc = 1,
+  omp_large_cap_mem_alloc = 2,
+  omp_const_mem_alloc = 3,
+  omp_high_bw_mem_alloc = 4,
+  omp_low_lat_mem_alloc = 5,
+  omp_cgroup_mem_alloc = 6,
+  omp_pteam_mem_alloc = 7,
+  omp_thread_mem_alloc = 8,
+  ompx_gnu_pinned_mem_alloc = 200,
+  ompx_gnu_managed_mem_alloc = 201,
+  __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+void f(int x)
+{
+  int a[x], *b[x];
+  int (&aRef)[x] = a;
+  int *(&bRef)[x] = b;
+
+  #pragma omp target private (aRef, bRef) \
+                    allocate (allocator (omp_low_lat_mem_alloc): aRef) \
+                    allocate (allocator (omp_large_cap_mem_alloc): bRef)
+  {
+    aRef[0] = 0;
+    bRef[0] = 0;
+  }
+}
+
+// Ensure that the size of memory allocated for the VLA is from a variable
+// rather than a constant, and that the default alignments are as expected.
+
+// { dg-final { scan-tree-dump "a\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "b\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target ilp32 } } }
+// { dg-final { scan-tree-dump "b\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 64\\\);" "omplower" { target { lp64 || llp64 } } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(4, D\\\.\[0-9\]\+, 5\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(4, D\\\.\[0-9\]\+, 2\\\);" "omplower" { target ilp32 } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(8, D\\\.\[0-9\]\+, 2\\\);" "omplower" { target { lp64 || llp64 } } } }
+// { dg-final { scan-tree-dump "aRef = D\\\.\[0-9\]\+;" "omplower" } }
+// { dg-final { scan-tree-dump "bRef = D\\\.\[0-9\]\+;" "omplower" } }

diff --git a/gcc/testsuite/g++.dg/gomp/pr113436-5.C b/gcc/testsuite/g++.dg/gomp/pr113436-5.C
new file mode 100644 (file)
index 0000000..f0a9533
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/pr113436-5.C
@@ -0,0 +1,28 @@
+// PR middle-end/113436
+// { dg-do "compile" }
+// { dg-options "-fopenmp -fdump-tree-omplower" }
+
+void f(int x)
+{
+  int a[x], *b[x];
+  int (&aRef)[x] = a;
+  int *(&bRef)[x] = b;
+
+  #pragma omp target private (aRef, bRef)
+  {
+    aRef[0] = 0;
+    bRef[0] = 0;
+  }
+}
+
+// Ensure that the size of memory allocated for the VLA is from a variable
+// rather than a constant, and that the default alignments are as expected.
+
+// { dg-final { scan-tree-dump "a\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "b\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target ilp32 } } }
+// { dg-final { scan-tree-dump "b\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 64\\\);" "omplower" { target { lp64 || llp64 } } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 32\\\);" "omplower" { target ilp32 } } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_alloca_with_align \\\(D\\\.\[0-9\]\+, 64\\\);" "omplower" { target { lp64 || llp64 } } } }
+// { dg-final { scan-tree-dump "aRef = D\\\.\[0-9\]\+;" "omplower" } }
+// { dg-final { scan-tree-dump "bRef = D\\\.\[0-9\]\+;" "omplower" } }

diff --git a/libgomp/testsuite/libgomp.c++/pr113436-1.C b/libgomp/testsuite/libgomp.c++/pr113436-1.C
index 0aae73b52cf851b5babdf5fdcd8f04fd364a0794..4e430f9c738678c175f1c4b4c71c05b9928adba5 100644 (file)
--- a/libgomp/testsuite/libgomp.c++/pr113436-1.C
+++ b/libgomp/testsuite/libgomp.c++/pr113436-1.C
@@ -21,7 +21,29 @@ test_int_by_ref ()
     }
 }
 
+void
+test_vla_by_ref (int n)
+{
+  int x[n];
+  for (int i = 0; i < n; i++)
+    x[i] = i;
+  int (&y)[n] = x;
+
+  #pragma omp target firstprivate(y) \
+                    allocate(allocator(omp_low_lat_mem_alloc), align(128): y)
+    {
+      if (((uintptr_t) &y) % 128  != 0)
+       __builtin_abort ();
+      for (int i = 0; i < n; i++)
+       y[i]++;
+      for (int i = 0; i < n; i++)
+       if (y[i] != i + 1)
+         __builtin_abort ();
+    }
+}
+
 int main ()
 {
   test_int_by_ref ();
+  test_vla_by_ref (17);
 }

diff --git a/libgomp/testsuite/libgomp.c++/pr113436-2.C b/libgomp/testsuite/libgomp.c++/pr113436-2.C
index 300399509897d11440eb9be7e158054fc9ce8f83..af54ca19df35277850a1d08d747f17acdf1b0283 100644 (file)
--- a/libgomp/testsuite/libgomp.c++/pr113436-2.C
+++ b/libgomp/testsuite/libgomp.c++/pr113436-2.C
@@ -19,7 +19,26 @@ test_int_by_ref ()
     }
 }
 
+void
+test_vla_by_ref (int n)
+{
+  int x[n];
+  for (int i = 0; i < n; i++)
+    x[i] = i;
+  int (&y)[n] = x;
+
+  #pragma omp target private(y) \
+                    allocate(allocator(omp_low_lat_mem_alloc), align(128): y)
+    {
+      if (((uintptr_t) &y) % 128  != 0)
+       __builtin_abort ();
+      for (int i = 0; i < n; i++)
+       y[i] = i + 1;
+    }
+}
+
 int main ()
 {
   test_int_by_ref ();
+  test_vla_by_ref (17);
 }