openmp: in_reduction support for Fortran

This patch implements support for the in_reduction clause for Fortran.
It also includes more completion of the taskgroup construct inside the
Fortran front-end, thus allowing task_reduction to work for task and
target constructs.

gcc/fortran/ChangeLog:

	* openmp.c (gfc_match_omp_clause_reduction): Add 'openmp_target' default
	false parameter. Add 'always,tofrom' map for OMP_LIST_IN_REDUCTION case.
	(gfc_match_omp_clauses): Add 'openmp_target' default false parameter,
	adjust call to gfc_match_omp_clause_reduction.
	(match_omp): Adjust call to gfc_match_omp_clauses
	* trans-openmp.c (gfc_trans_omp_taskgroup): Add call to
	gfc_match_omp_clause, create and return block.

gcc/ChangeLog:

	* omp-low.c (omp_copy_decl_2): For !ctx, use record_vars to add new copy
	as local variable.
	(scan_sharing_clauses): Place copy of OMP_CLAUSE_IN_REDUCTION decl in
	ctx->outer instead of ctx.

gcc/testsuite/ChangeLog:

	* gfortran.dg/gomp/reduction4.f90: Adjust omp target in_reduction' scan
	pattern.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/target-in-reduction-1.f90: New test.
	* testsuite/libgomp.fortran/target-in-reduction-2.f90: New test.

6 files changed, 147 insertions, 12 deletions

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 2a161f3..dcf22ac 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1202,7 +1202,7 @@ failed:
 
 static match
 gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
-				bool allow_derived)
+				bool allow_derived, bool openmp_target = false)
 {
   if (pc == 'r' && gfc_match ("reduction ( ") != MATCH_YES)
     return MATCH_NO;
@@ -1349,6 +1349,19 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
 	    n->u2.udr = gfc_get_omp_namelist_udr ();
 	    n->u2.udr->udr = udr;
 	  }
+	if (openmp_target && list_idx == OMP_LIST_IN_REDUCTION)
+	  {
+	    gfc_omp_namelist *p = gfc_get_omp_namelist (), **tl;
+	    p->sym = n->sym;
+	    p->where = p->where;
+	    p->u.map_op = OMP_MAP_ALWAYS_TOFROM;
+
+	    tl = &c->lists[OMP_LIST_MAP];
+	    while (*tl)
+	      tl = &((*tl)->next);
+	    *tl = p;
+	    p->next = NULL;
+	  }
      }
   return MATCH_YES;
 }
@@ -1417,7 +1430,8 @@ gfc_match_dupl_atomic (bool not_dupl, const char *name)
 static match
 gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		       bool first = true, bool needs_space = true,
-		       bool openacc = false, bool context_selector = false)
+		       bool openacc = false, bool context_selector = false,
+		       bool openmp_target = false)
 {
   bool error = false;
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
@@ -2121,8 +2135,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      goto error;
 	    }
 	  if ((mask & OMP_CLAUSE_IN_REDUCTION)
-	      && gfc_match_omp_clause_reduction (pc, c, openacc,
-						 allow_derived) == MATCH_YES)
+	      && gfc_match_omp_clause_reduction (pc, c, openacc, allow_derived,
+						 openmp_target) == MATCH_YES)
 	    continue;
 	  if ((mask & OMP_CLAUSE_INBRANCH)
 	      && (m = gfc_match_dupl_check (!c->inbranch && !c->notinbranch,
@@ -3578,7 +3592,8 @@ static match
 match_omp (gfc_exec_op op, const omp_mask mask)
 {
   gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, mask) != MATCH_YES)
+  if (gfc_match_omp_clauses (&c, mask, true, true, false, false,
+			     op == EXEC_OMP_TARGET) != MATCH_YES)
     return MATCH_ERROR;
   new_st.op = op;
   new_st.ext.omp_clauses = c;
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 4f1b346..aaeb950 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -6407,12 +6407,17 @@ gfc_trans_omp_task (gfc_code *code)
 static tree
 gfc_trans_omp_taskgroup (gfc_code *code)
 {
+  stmtblock_t block;
+  gfc_start_block (&block);
   tree body = gfc_trans_code (code->block->next);
   tree stmt = make_node (OMP_TASKGROUP);
   TREE_TYPE (stmt) = void_type_node;
   OMP_TASKGROUP_BODY (stmt) = body;
-  OMP_TASKGROUP_CLAUSES (stmt) = NULL_TREE;
-  return stmt;
+  OMP_TASKGROUP_CLAUSES (stmt) = gfc_trans_omp_clauses (&block,
+							code->ext.omp_clauses,
+							code->loc);
+  gfc_add_expr_to_block (&block, stmt);
+  return gfc_finish_block (&block);
 }
 
 static tree
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 057b7ae..15e4424 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -591,7 +591,15 @@ omp_copy_decl_2 (tree var, tree name, tree type, omp_context *ctx)
   tree copy = copy_var_decl (var, name, type);
 
   DECL_CONTEXT (copy) = current_function_decl;
-  DECL_CHAIN (copy) = ctx->block_vars;
+
+  if (ctx)
+    {
+      DECL_CHAIN (copy) = ctx->block_vars;
+      ctx->block_vars = copy;
+    }
+  else
+    record_vars (copy);
+
   /* If VAR is listed in task_shared_vars, it means it wasn't
      originally addressable and is just because task needs to take
      it's address.  But we don't need to take address of privatizations
@@ -602,7 +610,6 @@ omp_copy_decl_2 (tree var, tree name, tree type, omp_context *ctx)
 	  || (global_nonaddressable_vars
 	      && bitmap_bit_p (global_nonaddressable_vars, DECL_UID (var)))))
     TREE_ADDRESSABLE (copy) = 0;
-  ctx->block_vars = copy;
 
   return copy;
 }
@@ -1281,7 +1288,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		  tree at = t;
 		  if (ctx->outer)
 		    scan_omp_op (&at, ctx->outer);
-		  tree nt = omp_copy_decl_1 (at, ctx);
+		  tree nt = omp_copy_decl_1 (at, ctx->outer);
 		  splay_tree_insert (ctx->field_map,
 				     (splay_tree_key) &DECL_CONTEXT (t),
 				     (splay_tree_value) nt);
@@ -1322,7 +1329,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      tree at = decl;
 	      if (ctx->outer)
 		scan_omp_op (&at, ctx->outer);
-	      tree nt = omp_copy_decl_1 (at, ctx);
+	      tree nt = omp_copy_decl_1 (at, ctx->outer);
 	      splay_tree_insert (ctx->field_map,
 				 (splay_tree_key) &DECL_CONTEXT (decl),
 				 (splay_tree_value) nt);
diff --git a/gcc/testsuite/gfortran.dg/gomp/reduction4.f90 b/gcc/testsuite/gfortran.dg/gomp/reduction4.f90
index 52d504b..71b4231 100644
--- a/gcc/testsuite/gfortran.dg/gomp/reduction4.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/reduction4.f90
@@ -137,7 +137,7 @@ end
 ! { dg-final { scan-tree-dump-times "#pragma omp sections reduction\\(task,\\\+:a\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp simd linear\\(i:1\\) reduction\\(\\\+:a\\)" 2 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp simd linear\\(i:1\\) reduction\\(task,\\\+:a\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target in_reduction\\(\\\+:b\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(always,tofrom:b\\) in_reduction\\(\\\+:b\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp task in_reduction\\(\\\+:a\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp teams reduction\\(\\\+:b\\)" 2 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp taskloop reduction\\(\\\+:a\\) in_reduction\\(\\\+:b\\)" 2 "original" } }
diff --git a/libgomp/testsuite/libgomp.fortran/target-in-reduction-1.f90 b/libgomp/testsuite/libgomp.fortran/target-in-reduction-1.f90
new file mode 100644
index 0000000..f9acb71
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-in-reduction-1.f90
@@ -0,0 +1,78 @@
+! { dg-do run }
+
+module mod1
+  contains
+
+    subroutine foo (x, y)
+      integer :: x, y
+
+      !$omp taskgroup task_reduction (+: x, y)
+
+      !$omp target in_reduction (+: x, y)
+      x = x + 8
+      y = y + 16
+      !$omp end target
+
+      !$omp task in_reduction (+: x, y)
+      x = x + 2
+      y = y + 4
+      !$omp end task
+
+      !$omp end taskgroup
+    end subroutine foo
+
+    integer function bar (x)
+      integer, value :: x
+
+      !$omp taskgroup task_reduction (+: x)
+
+      !$omp target in_reduction (+: x)
+      x = x + 16
+      !$omp end target
+
+      !$omp task in_reduction (+: x)
+      x = x + 32
+      !$omp end task
+
+      !$omp end taskgroup
+
+      bar = x
+    end function bar
+  end module mod1
+
+program main
+  use mod1
+  integer :: x, y
+  real :: f;
+
+  x = 1
+  y = 1
+
+  call foo (x, y)
+
+  if (x .ne. 11) stop 1
+  if (y .ne. 21) stop 2
+
+  y = bar (8)
+  if (y .ne. 56) stop 3
+
+  x = 0
+  f = 0.0
+
+  !$omp taskgroup task_reduction (+: x, f)
+  !$omp target in_reduction (+: x, f)
+  x = x + 1
+  f = f + 2.0
+  !$omp end target
+
+  !$omp task in_reduction (+: x, f)
+  x = x + 2
+  f = f + 3.0
+  !$omp end task
+
+  !$omp end taskgroup
+
+  if (x .ne. 3) stop 4
+  if (f .ne. 5.0) stop 5
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target-in-reduction-2.f90 b/libgomp/testsuite/libgomp.fortran/target-in-reduction-2.f90
new file mode 100644
index 0000000..7f2e16b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-in-reduction-2.f90
@@ -0,0 +1,30 @@
+! { dg-do run }
+
+program main
+  integer :: x
+
+  x = 0
+  !$omp taskgroup task_reduction (+: x)
+  call foo (x)
+  call bar (x)
+  !$omp end taskgroup
+
+  if (x .ne. 3) stop 1
+
+contains
+
+  subroutine foo (x)
+    integer :: x
+    !$omp task in_reduction (+: x)
+    x = x + 1
+    !$omp end task
+  end subroutine foo
+
+  subroutine bar (x)
+    integer :: x
+    !$omp target in_reduction (+: x)
+    x = x + 2
+    !$omp end target
+  end subroutine bar
+
+end program main