openmp: Allocate memory for private/firstprivate clauses as directed by allocate...

This patch generates calls to GOMP_alloc to allocate memory for firstprivate
and private clauses on target constructs with an allocator and alignment
as specified by the allocate clause.

The decl values of the clause need to be adjusted to refer to the allocated
memory, and the initial values of variables need to be copied into the
allocated space for firstprivate variables.

For variable-length arrays, the size of the array is stored in a separate
variable, so the allocation and initialization need to be delayed until the
size is made available on the target.

gcc/

	PR middle-end/113436
	* omp-low.cc (is_variable_sized): Add extra is_ref argument.  Check
	referenced type if true.
	(lower_omp_target): Call lower_private_allocate to generate code to
	allocate memory for firstprivate/private clauses with allocators, and
	insert code after dependent variables have been initialized.
	Construct calls to free allocate memory and insert after target block.
	Adjust decl values for clause variables.  Copy value of firstprivate
	variables to allocated memory.

gcc/testsuite/

	PR middle-end/113436
	* c-c++-common/gomp/pr113436-1.c: New.
	* c-c++-common/gomp/pr113436-2.c: New.
	* g++.dg/gomp/pr113436.C: New.
	* gfortran.dg/gomp/pr113436-1.f90: New.
	* gfortran.dg/gomp/pr113436-2.f90: New.
	* gfortran.dg/gomp/pr113436-3.f90: New.
	* gfortran.dg/gomp/pr113436-4.f90: New.

libgomp/

	PR middle-end/113436
	* libgomp.texi (OpenMP 5.0): Mark allocate clause as implemented.
	(Memory allocation): Add documentation for use in target construct.
	* testsuite/libgomp.c++/firstprivate-1.C: Enable alignment check.
	* testsuite/libgomp.c++/pr113436-1.C: New.
	* testsuite/libgomp.c++/pr113436-2.C: New.
	* testsuite/libgomp.c++/private-1.C: Enable alignment check.
	* testsuite/libgomp.c-c++-common/pr113436-1.c: New.
	* testsuite/libgomp.c-c++-common/pr113436-2.c: New.
	* testsuite/libgomp.fortran/pr113436-1.f90: New.
	* testsuite/libgomp.fortran/pr113436-2.f90: New.

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 270a9abee3306f6c1a997ddedb01f4f5b0f4c636..15245eae9decb90ed7774cbcef3ba69054a34353 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -396,11 +396,15 @@ is_taskreg_ctx (omp_context *ctx)
   return is_parallel_ctx (ctx) || is_task_ctx (ctx) || is_host_teams_ctx (ctx);
 }
 
-/* Return true if EXPR is variable sized.  */
+/* Return true if EXPR is variable sized.  If IS_REF is true, then
+   EXPR is assumed to be a reference and the object it refers
+   to is checked instead.  */
 
 static inline bool
-is_variable_sized (const_tree expr)
+is_variable_sized (const_tree expr, bool is_ref = false)
 {
+  if (is_ref)
+    expr = TREE_TYPE (expr);
   return !TREE_CONSTANT (TYPE_SIZE_UNIT (TREE_TYPE (expr)));
 }
 
@@ -12813,10 +12817,17 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
   ilist = NULL;
   olist = NULL;
+
+  gimple_seq alloc_dlist = NULL;
+  hash_map<tree, tree> alloc_map;
+  hash_map<tree, gimple_seq> alloc_seq_map;
+
   for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
     switch (OMP_CLAUSE_CODE (c))
       {
-       tree var, x;
+       tree var, x, new_var, allocator, allocate_ptr, size;
+       gimple_seq alloc_seq;
+       bool by_ref;
 
       default:
        break;
@@ -12997,10 +13008,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
          }
        map_cnt++;
        var = OMP_CLAUSE_DECL (c);
-       if (!omp_privatize_by_reference (var)
-           && !is_gimple_reg_type (TREE_TYPE (var)))
+       new_var = lookup_decl (var, ctx);
+       allocator = NULL_TREE;
+       allocate_ptr = NULL_TREE;
+       size = TREE_TYPE (var);
+       by_ref = omp_privatize_by_reference (var);
+       if (by_ref)
+         size = TREE_TYPE (size);
+       size = TYPE_SIZE_UNIT (size);
+       if (is_variable_sized (var, by_ref))
+         size = lookup_decl (size, ctx);
+       alloc_seq = NULL;
+       if (lower_private_allocate (var, new_var, allocator, allocate_ptr,
+                                   &alloc_seq, ctx, by_ref, size))
+         {
+           alloc_map.put (new_var, allocate_ptr);
+           alloc_seq_map.put (new_var, alloc_seq);
+         }
+       if (!by_ref && !is_gimple_reg_type (TREE_TYPE (var)))
          {
-           tree new_var = lookup_decl (var, ctx);
            if (is_variable_sized (var))
              {
                tree pvar = DECL_VALUE_EXPR (var);
@@ -13011,6 +13037,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                x = build_fold_indirect_ref (new_pvar);
                TREE_THIS_NOTRAP (x) = 1;
              }
+           else if (allocate_ptr)
+             x = build_fold_indirect_ref (allocate_ptr);
            else
              x = build_receiver_ref (var, true, ctx);
            SET_DECL_VALUE_EXPR (new_var, x);
@@ -13020,6 +13048,26 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
          if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR
              && lang_hooks.decls.omp_array_data (var, true))
            map_cnt += 2;
+
+      do_dtor:
+       if (allocator)
+         {
+           if (!is_gimple_val (allocator))
+             {
+               tree avar = create_tmp_var (TREE_TYPE (allocator));
+               gimplify_assign (avar, allocator, &alloc_dlist);
+               allocator = avar;
+             }
+           if (!is_gimple_val (allocate_ptr))
+             {
+               tree apvar = create_tmp_var (TREE_TYPE (allocate_ptr));
+               gimplify_assign (apvar, allocate_ptr, &alloc_dlist);
+               allocate_ptr = apvar;
+             }
+           tree f = builtin_decl_explicit (BUILT_IN_GOMP_FREE);
+           gimple *g = gimple_build_call (f, 2, allocate_ptr, allocator);
+           gimple_seq_add_stmt (&alloc_dlist, g);
+         }
        break;
 
       case OMP_CLAUSE_PRIVATE:
@@ -13034,7 +13082,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            break;
          }
        var = OMP_CLAUSE_DECL (c);
-       if (is_variable_sized (var))
+       new_var = lookup_decl (var, ctx);
+       allocator = NULL_TREE;
+       allocate_ptr = NULL_TREE;
+       alloc_seq = NULL;
+       size = TREE_TYPE (var);
+       by_ref = omp_privatize_by_reference (var);
+       if (by_ref)
+         size = TREE_TYPE (size);
+       size = TYPE_SIZE_UNIT (size);
+       if (is_variable_sized (var, by_ref))
+         size = lookup_decl (size, ctx);
+       lower_private_allocate (var, new_var, allocator, allocate_ptr,
+                               &alloc_seq, ctx, by_ref, size);
+       if (allocate_ptr)
+         {
+           alloc_map.put (new_var, allocate_ptr);
+           alloc_seq_map.put (new_var, alloc_seq);
+         }
+       if (!allocate_ptr && is_variable_sized (var))
          {
            tree new_var = lookup_decl (var, ctx);
            tree pvar = DECL_VALUE_EXPR (var);
@@ -13047,7 +13113,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            SET_DECL_VALUE_EXPR (new_var, x);
            DECL_HAS_VALUE_EXPR_P (new_var) = 1;
          }
-       break;
+       goto do_dtor;
 
       case OMP_CLAUSE_USE_DEVICE_PTR:
       case OMP_CLAUSE_USE_DEVICE_ADDR:
@@ -13961,10 +14027,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   if (offloaded || data_region)
     {
       tree prev = NULL_TREE;
+      bool by_ref;
       for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
        switch (OMP_CLAUSE_CODE (c))
          {
-           tree var, x;
+           tree var, x, new_var, *allocate_ptr;
          default:
            break;
          case OMP_CLAUSE_FIRSTPRIVATE:
@@ -13972,13 +14039,34 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            if (is_gimple_omp_oacc (ctx->stmt))
              break;
            var = OMP_CLAUSE_DECL (c);
-           if (omp_privatize_by_reference (var)
-               || is_gimple_reg_type (TREE_TYPE (var)))
+           new_var = lookup_decl (var, ctx);
+           allocate_ptr = alloc_map.get (new_var);
+           by_ref = omp_privatize_by_reference (var);
+           if (allocate_ptr)
+             {
+               if (is_variable_sized (var, by_ref))
+                 /* Handle this in the next pass when the size is
+                    available.  */
+                 break;
+
+               gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+               gcc_assert (allocate_seq);
+               gimple_seq_add_seq (&new_body, *allocate_seq);
+
+               if (by_ref)
+                 {
+                   x = fold_convert (TREE_TYPE (new_var), *allocate_ptr);
+                   gimplify_assign (new_var, x, &new_body);
+                   new_var = build_fold_indirect_ref (new_var);
+                 }
+               else
+                 new_var = build_simple_mem_ref (*allocate_ptr);
+             }
+           if (by_ref || is_gimple_reg_type (TREE_TYPE (var)))
              {
-               tree new_var = lookup_decl (var, ctx);
                tree type;
                type = TREE_TYPE (var);
-               if (omp_privatize_by_reference (var))
+               if (by_ref)
                  type = TREE_TYPE (type);
                if ((INTEGRAL_TYPE_P (type)
                     && TYPE_PRECISION (type) <= POINTER_SIZE)
@@ -13990,7 +14078,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                    x = fold_convert (type, x);
                    gimplify_expr (&x, &new_body, NULL, is_gimple_val,
                                   fb_rvalue);
-                   if (omp_privatize_by_reference (var))
+                   if (by_ref && !allocate_ptr)
                      {
                        tree v = create_tmp_var_raw (type, get_name (var));
                        gimple_add_tmp_var (v);
@@ -13999,17 +14087,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                                             gimple_build_assign (v, x));
                        x = build_fold_addr_expr (v);
                      }
-                   gimple_seq_add_stmt (&new_body,
-                                        gimple_build_assign (new_var, x));
+                   gimplify_assign (new_var, x, &new_body);
                  }
                else
                  {
-                   bool by_ref = !omp_privatize_by_reference (var);
-                   x = build_receiver_ref (var, by_ref, ctx);
+                   x = build_receiver_ref (var, allocate_ptr || !by_ref, ctx);
                    gimplify_expr (&x, &new_body, NULL, is_gimple_val,
                                   fb_rvalue);
-                   gimple_seq_add_stmt (&new_body,
-                                        gimple_build_assign (new_var, x));
+                   gimplify_assign (new_var, x, &new_body);
                  }
              }
            else if (is_variable_sized (var))
@@ -14024,29 +14109,59 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                gimple_seq_add_stmt (&new_body,
                                     gimple_build_assign (new_var, x));
              }
+           else if (allocate_ptr)
+             {
+               x = build_receiver_ref (var, true, ctx);
+               new_var = unshare_expr (new_var);
+               x = lang_hooks.decls.omp_clause_copy_ctor (c, new_var, x);
+               gimplify_and_add (x, &new_body);
+             }
            break;
          case OMP_CLAUSE_PRIVATE:
            if (is_gimple_omp_oacc (ctx->stmt))
              break;
            var = OMP_CLAUSE_DECL (c);
-           if (omp_privatize_by_reference (var))
+           new_var = lookup_decl (var, ctx);
+           allocate_ptr = alloc_map.get (new_var);
+           by_ref = omp_privatize_by_reference (var);
+           if (allocate_ptr)
+             {
+               if (is_variable_sized (var, by_ref))
+                 /* Handle this in the next pass when the size is
+                    available.  */
+                 break;
+
+               gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+               gcc_assert (allocate_seq);
+               gimple_seq_add_seq (&new_body, *allocate_seq);
+
+               if (!by_ref)
+                 new_var = build_simple_mem_ref (*allocate_ptr);
+             }
+           if (by_ref)
              {
                location_t clause_loc = OMP_CLAUSE_LOCATION (c);
-               tree new_var = lookup_decl (var, ctx);
-               x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
-               if (TREE_CONSTANT (x))
+               if (!allocate_ptr)
                  {
-                   x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
-                                           get_name (var));
-                   gimple_add_tmp_var (x);
-                   TREE_ADDRESSABLE (x) = 1;
-                   x = build_fold_addr_expr_loc (clause_loc, x);
+                   x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
+                   if (TREE_CONSTANT (x))
+                     {
+                       x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
+                                               get_name (var));
+                       gimple_add_tmp_var (x);
+                       TREE_ADDRESSABLE (x) = 1;
+                       x = build_fold_addr_expr_loc (clause_loc, x);
+                     }
+                   else
+                     break;
+
+                   x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
+                   gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+                                  fb_rvalue);
                  }
                else
-                 break;
+                 x = *allocate_ptr;
 
-               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));
              }
@@ -14055,7 +14170,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
          case OMP_CLAUSE_USE_DEVICE_ADDR:
          case OMP_CLAUSE_HAS_DEVICE_ADDR:
          case OMP_CLAUSE_IS_DEVICE_PTR:
-           tree new_var;
            gimple_seq assign_body;
            bool is_array_data;
            bool do_optional_check;
@@ -14343,26 +14457,43 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            break;
          case OMP_CLAUSE_PRIVATE:
            var = OMP_CLAUSE_DECL (c);
-           if (is_variable_sized (var))
+           by_ref = omp_privatize_by_reference (var);
+           if (is_variable_sized (var, by_ref))
              {
-               location_t clause_loc = OMP_CLAUSE_LOCATION (c);
                tree new_var = lookup_decl (var, ctx);
-               tree pvar = DECL_VALUE_EXPR (var);
-               gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
-               pvar = TREE_OPERAND (pvar, 0);
+               tree *allocate_ptr = alloc_map.get (new_var);
+               if (allocate_ptr)
+                 {
+                   gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+                   gcc_assert (allocate_seq);
+                   gimple_seq_add_seq (&new_body, *allocate_seq);
+                 }
+               location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+               tree pvar = var;
+               if (!by_ref)
+                 {
+                   pvar = DECL_VALUE_EXPR (var);
+                   gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+                   pvar = TREE_OPERAND (pvar, 0);
+                 }
                gcc_assert (DECL_P (pvar));
                tree new_pvar = lookup_decl (pvar, ctx);
-               tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
-               tree al = size_int (DECL_ALIGN (var));
-               tree x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
-               x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
-               x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
+               tree x;
+               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));
+                   x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
+                   x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
+                 }
+               else
+                 x = *allocate_ptr;
                gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
                gimple_seq_add_stmt (&new_body,
                                     gimple_build_assign (new_pvar, x));
              }
-           else if (omp_privatize_by_reference (var)
-                    && !is_gimple_omp_oacc (ctx->stmt))
+           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);
@@ -14384,6 +14515,42 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                                     gimple_build_assign (new_var, x));
              }
            break;
+         case OMP_CLAUSE_FIRSTPRIVATE:
+           var = OMP_CLAUSE_DECL (c);
+           by_ref = omp_privatize_by_reference (var);
+           if (is_variable_sized (var, by_ref))
+             {
+               tree new_var = lookup_decl (var, ctx);
+               tree *allocate_ptr = alloc_map.get (new_var);
+               if (!allocate_ptr)
+                 break;
+               gimple_seq *allocate_seq = alloc_seq_map.get (new_var);
+               gcc_assert (allocate_seq);
+               gimple_seq_add_seq (&new_body, *allocate_seq);
+
+               location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+               tree pvar = var;
+               if (!by_ref)
+                 {
+                   pvar = DECL_VALUE_EXPR (var);
+                   gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+                   pvar = TREE_OPERAND (pvar, 0);
+                 }
+               gcc_assert (DECL_P (pvar));
+               tree new_pvar = lookup_decl (pvar, ctx);
+               tree x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar),
+                                          *allocate_ptr);
+               gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+               gimple_seq_add_stmt (&new_body,
+                                    gimple_build_assign (new_pvar, x));
+
+               x = build_receiver_ref (var, true, ctx);
+               new_var = unshare_expr (new_var);
+               if (by_ref)
+                 new_var = build_fold_indirect_ref (new_var);
+               x = lang_hooks.decls.omp_clause_copy_ctor (c, new_var, x);
+               gimplify_and_add (x, &new_body);
+             }
          }
 
       gimple_seq fork_seq = NULL;
@@ -14408,6 +14575,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       gimple_seq_add_seq (&new_body, fork_seq);
       gimple_seq_add_seq (&new_body, tgt_body);
       gimple_seq_add_seq (&new_body, join_seq);
+      gimple_seq_add_seq (&new_body, alloc_dlist);
 
       if (offloaded)
        {

diff --git a/gcc/testsuite/c-c++-common/gomp/pr113436-1.c b/gcc/testsuite/c-c++-common/gomp/pr113436-1.c
new file mode 100644 (file)
index 0000000..985cc21
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-1.c
@@ -0,0 +1,31 @@
+/* PR middle-end/113436 */
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#include <omp.h>
+
+void
+f()
+{
+  int A, B[10], *C;
+  A = 5;
+  C = (int *) __builtin_malloc (sizeof (int) * 10);
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+      
+  #pragma omp target private(A) private(B) private(C) allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+    {
+      A = 99;
+      for (int i = 0; i < 10; i++)
+        B[i] = -i-23;
+      C = &A;
+    }
+}
+
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128, 4, 5\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128, 40, 5\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128, 8, 5\\\);" "omplower" { target { lp64 || llp64 } } } } */
+/* { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[i\\\] = D\\\.\[0-9\]\+;" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = D\\\.\[0-9\]\+;" "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]\+, 5\\\);" 3 "omplower" } } */

diff --git a/gcc/testsuite/c-c++-common/gomp/pr113436-2.c b/gcc/testsuite/c-c++-common/gomp/pr113436-2.c
new file mode 100644 (file)
index 0000000..1755b6b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr113436-2.c
@@ -0,0 +1,32 @@
+/* PR middle-end/113436 */
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#include <omp.h>
+
+void
+g()
+{
+  int A, B[10], *C;
+  A = 5;
+  C = (int *) __builtin_malloc (sizeof (int) * 10);
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+      
+  #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) allocate(allocator(omp_high_bw_mem_alloc), align(64): A, B, C)
+    {
+      A = 99;
+      for (int i = 0; i < 10; i++)
+        B[i] = -i-23;
+      C = &A;
+    }
+}
+
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, 4, 4\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, 40, 4\\\);" "omplower" { target int32 } } } */
+/* { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, 8, 4\\\);" "omplower" { target { lp64 || llp64 } } } } */
+/* { dg-final { scan-tree-dump-times "\\\*D\\\.\[0-9\]\+ = D\\\.\[0-9\]\+;" 3 "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\) = \\\(\\\*D\\\.\[0-9\]\+\\\);" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } } */
+/* { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[i\\\] = D\\\.\[0-9\]\+;" "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]+, 4\\\)" 3 "omplower" } } */

diff --git a/gcc/testsuite/g++.dg/gomp/pr113436.C b/gcc/testsuite/g++.dg/gomp/pr113436.C
new file mode 100644 (file)
index 0000000..ad1cc2f
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/pr113436.C
@@ -0,0 +1,22 @@
+// PR middle-end/113436
+// { dg-do "compile" }
+// { dg-options "-std=gnu++20 -fopenmp -fdump-tree-omplower" }
+
+#include <omp.h>
+
+void f()
+{
+  int a[10];
+  auto &aRef = a;
+
+  #pragma omp target firstprivate(aRef) \
+                    allocate(align(128), allocator(omp_low_lat_mem_alloc): aRef)
+    aRef[0] = 1;
+}
+
+// { dg-final { scan-tree-dump "aRef\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(128, 40, 5\\\);" "omplower" { target int32 } } }
+// { dg-final { scan-tree-dump "aRef = aRef\\\.\[0-9\]\+;" "omplower" } }
+// { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = \\\(\\\*D\\\.\[0-9\]+\\\);" "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 \\\(aRef\\\.\[0-9\]\+, 5\\\);" "omplower" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90 b/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
new file mode 100644 (file)
index 0000000..da757ed
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-1.f90
@@ -0,0 +1,35 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+program g
+  use omp_lib
+  implicit none
+
+  integer :: A, B(10)
+  integer, allocatable :: C(:)
+  integer :: i
+
+  A = 5;
+  allocate(C(10))
+  do i = 1, 10
+    B(i) = i + 5
+    C(i) = B(i)
+  end do
+      
+  !$omp target private(A) private(B) private(C) allocate(allocator(omp_high_bw_mem_alloc), align(64): A, B, C)
+      A = 99
+      do i = 1, 10
+        B(i) = -i - 23
+        C(i) = i + 23
+      end do
+  !$omp end target
+end program g
+
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, 4, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, 40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump-times "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, \[0-9\]\+, 4\\\);" 3 "omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+\\\]\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]+, 4\\\)" 3 "omplower" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90
new file mode 100644 (file)
index 0000000..0eaf8b5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-2.f90
@@ -0,0 +1,38 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+program g
+  use omp_lib
+  implicit none
+
+  integer :: A, B(10)
+  integer, allocatable :: C(:)
+  integer :: i
+
+  A = 5;
+  allocate(C(10))
+  do i = 1, 10
+    B(i) = i + 5
+    C(i) = B(i)
+  end do
+      
+  !$omp target firstprivate(A) firstprivate(B) firstprivate(C) allocate(allocator(omp_high_bw_mem_alloc), align(64): A, B, C)
+    A = 99
+    do i = 1, 10
+      B(i) = -i - 23
+      C(i) = i + 23
+    end do
+  !$omp end target
+end program g
+
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, 4, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, 40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump-times "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(64, \[0-9\]\+, 4\\\);" 3 "omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\) = \\\(\\\*D\\\.\[0-9\]\+\\\);" "omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = \\\*D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump "\\\*D\\\.\[0-9\]\+ = 99;" "omplower" } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+\\\]\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\\(D\\\.\[0-9\]+, 4\\\)" 3 "omplower" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90 b/gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90
new file mode 100644 (file)
index 0000000..f04200f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-3.f90
@@ -0,0 +1,26 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fno-automatic -fdump-tree-omplower" }
+
+program g
+  use omp_lib
+  implicit none
+
+  integer :: A(10)
+  integer :: i
+
+  do i = 1, 10
+    A(i) = i + 5
+  end do
+      
+  !$omp target private(A) allocate(allocator(omp_high_bw_mem_alloc), align(16): A)
+      do i = 1, 10
+        A(i) = -i + 23
+      end do
+  !$omp end target
+end program g
+
+! { dg-excess-errors "Flag '-fno-automatic' overwrites '-frecursive' implied by '-fopenmp'" }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(16, 40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_free \\\(D\\\.\[0-9\]+, 4\\\)" "omplower" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90 b/gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90
new file mode 100644 (file)
index 0000000..19889d2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr113436-4.f90
@@ -0,0 +1,27 @@
+! PR middle-end/113436
+! { dg-do compile }
+! { dg-options "-fopenmp -fno-automatic -fdump-tree-omplower" }
+
+program g
+  use omp_lib
+  implicit none
+
+  integer :: A(10)
+  integer :: i
+
+  do i = 1, 10
+    A(i) = i + 5
+  end do
+      
+  !$omp target firstprivate(A) allocate(allocator(omp_high_bw_mem_alloc), align(32): A)
+      do i = 1, 10
+        A(i) = -i + 23
+      end do
+  !$omp end target
+end program g
+
+! { dg-excess-errors "Flag '-fno-automatic' overwrites '-frecursive' implied by '-fopenmp'" }
+! { dg-final { scan-tree-dump "D\\\.\[0-9\]\+ = __builtin_GOMP_alloc \\\(32, 40, 4\\\);" "omplower" { target int32 } } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\) = \\\(\\\*D\\\.\[0-9\]\+\\\);" "omplower" } }
+! { dg-final { scan-tree-dump "\\\(\\\*D\\\.\[0-9\]\+\\\)\\\[D\\\.\[0-9\]\+\\\] = D\\\.\[0-9\]\+;" "omplower" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_free \\\(D\\\.\[0-9\]+, 4\\\)" "omplower" } }

diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index eab41ac3508a3c7fc9c2e8d2cc86086125caeeac..258ea8a76198ef3a6a142d62e797feac7ffbe1a0 100644 (file)
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -234,8 +234,7 @@ The OpenMP 4.5 specification is fully supported.
 @item Memory management routines @tab Y @tab
 @item @code{allocate} directive @tab P
       @tab C++ unsupported; see also @ref{Memory allocation}
-@item @code{allocate} clause @tab P @tab Clause has no effect on @code{target}
-      (@uref{https://gcc.gnu.org/PR113436,PR113436})
+@item @code{allocate} clause @tab Y @tab
 @item @code{use_device_addr} clause on @code{target data} @tab Y @tab
 @item @code{ancestor} modifier on @code{device} clause @tab Y @tab
 @item Implicit declare target directive @tab Y @tab
@@ -6895,6 +6894,11 @@ The description below applies to:
       constant expression with value @code{omp_default_mem_alloc} and no
       @code{align} modifier has been specified. (In that case, the normal
       @code{malloc} allocation is used.)
+@item The @code{allocate} clause can be used in the @code{target} construct
+      to specify the memory used by @code{private} and @code{firstprivate}
+      variables on offload devices.  In the case of @code{firstprivate}, the
+      initial data is first allocated using the default memory allocator, then
+      copied to the memory region specified by the allocator.
 @item The @code{allocate} directive for variables in static memory; while
       the alignment is honored, the normal static memory is used.
 @item Using the @code{allocate} directive for automatic/stack variables, except

diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-1.C b/libgomp/testsuite/libgomp.c++/firstprivate-1.C
index ae5d4fbe1bf9323bd3dc2064bc0db6e11c679e95..a7393382cf046c313aae85ff9d5d21a648d9ec3b 100644 (file)
--- a/libgomp/testsuite/libgomp.c++/firstprivate-1.C
+++ b/libgomp/testsuite/libgomp.c++/firstprivate-1.C
@@ -90,14 +90,13 @@ S::g (int dev)
                       allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
                       device(dev)
     {
-#if 0  /* FIXME: The following is disabled because of PR middle-end/113436.  */
       if (((uintptr_t) &A) % 128  != 0)
        abort ();
       if (((uintptr_t) &B) % 128  != 0)
        abort ();
       if (((uintptr_t) &C) % 128  != 0)
        abort ();
-#endif
+
       if (A != 5)
        abort ();
       for (int i = 0; i < 10; i++)
@@ -227,14 +226,13 @@ St<T>::gt (int dev)
                      allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
                      device(dev)
     {
-#if 0  /* FIXME: The following is disabled because of PR middle-end/113436.  */
       if (((uintptr_t) &A) % 128  != 0)
        abort ();
       if (((uintptr_t) &B) % 128  != 0)
        abort ();
       if (((uintptr_t) &C) % 128  != 0)
        abort ();
-#endif
+
       if (A != 5)
        abort ();
       for (int i = 0; i < 10; i++)

diff --git a/libgomp/testsuite/libgomp.c++/pr113436-1.C b/libgomp/testsuite/libgomp.c++/pr113436-1.C
new file mode 100644 (file)
index 0000000..0aae73b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr113436-1.C
@@ -0,0 +1,27 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_ref ()
+{
+  int a = 5;
+  int &b = a;
+
+  #pragma omp target firstprivate(b) \
+                    allocate(allocator(omp_high_bw_mem_alloc), align(64): b)
+    {
+      if (((uintptr_t) &b) % 64  != 0)
+       __builtin_abort ();
+      b *= 7;
+      if (b != 35)
+       __builtin_abort ();
+    }
+}
+
+int main ()
+{
+  test_int_by_ref ();
+}

diff --git a/libgomp/testsuite/libgomp.c++/pr113436-2.C b/libgomp/testsuite/libgomp.c++/pr113436-2.C
new file mode 100644 (file)
index 0000000..3003995
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr113436-2.C
@@ -0,0 +1,25 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_ref ()
+{
+  int a = 5;
+  int &b = a;
+
+  #pragma omp target private(b) \
+                    allocate(allocator(omp_high_bw_mem_alloc), align(64): b)
+    {
+      if (((uintptr_t) &b) % 64  != 0)
+       __builtin_abort ();
+      b = 7;
+    }
+}
+
+int main ()
+{
+  test_int_by_ref ();
+}

diff --git a/libgomp/testsuite/libgomp.c++/private-1.C b/libgomp/testsuite/libgomp.c++/private-1.C
index 19ee726a2222552404b9ffbf638aa885c80982ab..84bfc8225f17a31dd0d551e6321fab08f9d19700 100644 (file)
--- a/libgomp/testsuite/libgomp.c++/private-1.C
+++ b/libgomp/testsuite/libgomp.c++/private-1.C
@@ -75,14 +75,13 @@ S::g (int dev)
                      allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
                      device(dev)
     {
-#if 0  /* FIXME: The following is disabled because of PR middle-end/113436.  */
       if (((uintptr_t) &A) % 128  != 0)
        abort ();
       if (((uintptr_t) &B) % 128  != 0)
        abort ();
       if (((uintptr_t) &C) % 128  != 0)
        abort ();
-#endif
+
       A = 99;
       for (int i = 0; i < 10; i++)
        B[i] = -i-23;

diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr113436-1.c b/libgomp/testsuite/libgomp.c-c++-common/pr113436-1.c
new file mode 100644 (file)
index 0000000..18a8792
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr113436-1.c
@@ -0,0 +1,94 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_val ()
+{
+  int x = 64;
+
+  #pragma omp target firstprivate(x) \
+                    allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+    {
+      if (((uintptr_t) &x) % 16  != 0)
+       __builtin_abort ();
+      x *= 2;
+      if (x != 128)
+       __builtin_abort ();
+    }
+}
+
+void
+test_struct_by_val ()
+{
+  struct S {
+    int a[4];
+    float b[4];
+  } s = { { 1, 2, 3, 4 }, { 5.0f, 6.0f, 7.0f, 8.0f } };
+
+  #pragma omp target firstprivate(s) \
+                    allocate(allocator(omp_low_lat_mem_alloc), align(32): s)
+    {
+      if (((uintptr_t) &s) % 32  != 0)
+       __builtin_abort ();
+      for (int i = 0; i < 4; i++)
+       {
+         s.a[i] *= 2;
+         s.b[i] *= 2.0f;
+       }
+      for (int i = 0; i < 4; i++)
+       if (s.a[i] != (i + 1) * 2 || s.b[i] != (i + 5) * 2.0f)
+         __builtin_abort ();
+    }
+}
+
+void
+test_ptr ()
+{
+  int x = 42;
+  int *p = &x;
+  uintptr_t p_orig = (uintptr_t) p;
+  uintptr_t p_new;
+
+  #pragma omp target firstprivate(p) \
+                    allocate(allocator(omp_default_mem_alloc), align(16): p) \
+                    map(from: p_new)
+    {
+      if (((uintptr_t) &p) % 16  != 0)
+       __builtin_abort ();
+      p_new = (uintptr_t) p;
+    }
+
+  if (p_new != p_orig)
+      __builtin_abort ();
+}
+
+void
+test_vla (int n)
+{
+  int x[n];
+  for (int i = 0; i < n; i++)
+    x[i] = i;
+
+  #pragma omp target firstprivate(x) \
+                    allocate(allocator(omp_high_bw_mem_alloc), align(128): x)
+    {
+      if (((uintptr_t) &x) % 128  != 0)
+       __builtin_abort ();
+      for (int i = 0; i < n; i++)
+       x[i]++;
+      for (int i = 0; i < n; i++)
+       if (x[i] != i + 1)
+         __builtin_abort ();
+    }
+}
+
+int main ()
+{
+  test_int_by_val ();
+  test_struct_by_val ();
+  test_ptr ();
+  test_vla (16);
+}

diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr113436-2.c b/libgomp/testsuite/libgomp.c-c++-common/pr113436-2.c
new file mode 100644 (file)
index 0000000..117944a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr113436-2.c
@@ -0,0 +1,80 @@
+/* PR middle-end/113436 */
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdint.h>
+
+void
+test_int_by_val ()
+{
+  int x;
+
+  #pragma omp target private(x) \
+                    allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+    {
+      if (((uintptr_t) &x) % 16  != 0)
+       __builtin_abort ();
+      x = 2;
+    }
+}
+
+void
+test_struct_by_val ()
+{
+  struct S {
+    int a[4];
+    float b[4];
+  } s = { { 1, 2, 3, 4 }, { 5.0f, 6.0f, 7.0f, 8.0f } };
+
+  #pragma omp target private(s) \
+                    allocate(allocator(omp_low_lat_mem_alloc), align(32): s)
+    {
+      if (((uintptr_t) &s) % 32  != 0)
+       __builtin_abort ();
+      for (int i = 0; i < 4; i++)
+       {
+         s.a[i] = i + 1;
+         s.b[i] = 2.0f * i;
+       }
+    }
+}
+
+void
+test_ptr ()
+{
+  int x = 42;
+  int *p = &x;
+
+  #pragma omp target firstprivate(p) \
+                    allocate(allocator(omp_default_mem_alloc), align(16): p)
+    {
+      if (((uintptr_t) &p) % 16  != 0)
+       __builtin_abort ();
+      p++;
+    }
+}
+
+void
+test_vla (int n)
+{
+  int x[n];
+  for (int i = 0; i < n; i++)
+    x[i] = i;
+
+  #pragma omp target private(x) \
+                    allocate(allocator(omp_high_bw_mem_alloc), align(128): x)
+    {
+      if (((uintptr_t) &x) % 128  != 0)
+       __builtin_abort ();
+      for (int i = 0; i < n; i++)
+       x[i] = i * 2;
+    }
+}
+
+int main ()
+{
+  test_int_by_val ();
+  test_struct_by_val ();
+  test_ptr ();
+  test_vla (32);
+}

diff --git a/libgomp/testsuite/libgomp.fortran/pr113436-1.f90 b/libgomp/testsuite/libgomp.fortran/pr113436-1.f90
new file mode 100644 (file)
index 0000000..0251525
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr113436-1.f90
@@ -0,0 +1,67 @@
+! PR middle-end/113436
+! { dg-do run }
+
+program main
+  use omp_lib
+  implicit none
+
+  call test_integer
+  call test_derived_type
+  call test_vla
+contains
+  subroutine test_integer
+    integer :: x = 64
+
+    !$omp target firstprivate(x) &
+    !$omp & allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+      if (mod (loc (x), 16) /= 0) stop 1
+      x = x * 2
+      if (x /= 128) stop 2
+    !$omp end target
+  end subroutine
+
+  subroutine test_derived_type
+    type :: Ty
+      integer :: a(4)
+      real*4 :: b(4)
+    end type
+    type (Ty) :: t = Ty (a=(/1, 2, 3, 4/), b=(/5.0, 6.0, 7.0, 8.0/))
+    integer :: i
+
+    !$omp target firstprivate(t) &
+    !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(32): t)
+      if (mod (loc (t), 32) /= 0) stop 3
+      do i = 1, 4
+       t%a(i) = t%a(i) * 2
+       t%b(i) = t%b(i) * 2.0
+      end do
+      do i = 1, 4
+       if (t%a(i) /= i * 2) stop 4
+       if (t%b(i) /= (i + 4) * 2.0) stop 5
+      end do
+    !$omp end target
+  end subroutine
+
+  subroutine test_vla
+    integer :: n = 10
+    integer :: i
+    block
+      integer :: a(n)
+
+      do i = 1, n
+       a(i) = i * 3
+      end do
+
+      !$omp target firstprivate(a) &
+      !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(64): a)
+        if (mod (loc (a), 64) /= 0) stop 6
+        do i = 1, n
+         a(i) = a(i) * 2
+        end do
+        do i = 1, n
+         if (a(i) /= i * 6) stop 7
+        end do
+      !$omp end target
+    end block
+  end subroutine
+end program

diff --git a/libgomp/testsuite/libgomp.fortran/pr113436-2.f90 b/libgomp/testsuite/libgomp.fortran/pr113436-2.f90
new file mode 100644 (file)
index 0000000..2ab257b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr113436-2.f90
@@ -0,0 +1,58 @@
+! PR middle-end/113436
+! { dg-do run }
+
+program main
+  use omp_lib
+  implicit none
+
+  call test_integer
+  call test_derived_type
+contains
+  subroutine test_integer
+    integer :: x
+
+    !$omp target private(x) &
+    !$omp & allocate(allocator(omp_high_bw_mem_alloc), align(16): x)
+      if (mod (loc (x), 16) /= 0) stop 1
+      x = 2
+    !$omp end target
+  end subroutine
+
+  subroutine test_derived_type
+    type :: Ty
+      integer :: a(4)
+      real*4 :: b(4)
+    end type
+    type (Ty) :: t
+    integer :: i
+
+    !$omp target private(t) &
+    !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(32): t)
+      if (mod (loc (t), 32) /= 0) stop 2
+      do i = 1, 4
+       t%a(i) = i
+       t%b(i) = i * 2.0
+      end do
+    !$omp end target
+  end subroutine
+
+  subroutine test_vla
+    integer :: n = 10
+    integer :: i
+    block
+      integer :: a(n)
+
+      do i = 1, n
+       a(i) = i * 3
+      end do
+
+      !$omp target firstprivate(a) &
+      !$omp & allocate(allocator(omp_low_lat_mem_alloc), align(64): a)
+        if (mod (loc (a), 64) /= 0) stop 6
+        do i = 1, n
+         a(i) = a(i) * 2
+        end do
+      !$omp end target
+    end block
+  end subroutine
+end program