Hi Jakub, and Fortran folks,
this patch does the required adjustments to let 'in_reduction' work for Fortran.
Not just for the target directive actually, task directive is also working after
this patch.
There is a little bit of adjustment in omp-low.c:scan_sharing_clauses:
RTL expand of the copy of the OMP_CLAUSE_IN_REDUCTION decl was failing
for Fortran by-reference arguments, which seems to work after placing them
under the outer ctx (when it exists). This also now needs checking the field_map
for existence of the field before inserting.
Tested without regressions on mainline trunk, is this okay?
(testing for devel/omp/gcc-11 is in progress)
Thanks,
Chung-Lin
2021-09-17 Chung-Lin Tang <[email protected]>
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 (scan_sharing_clauses): Place in_reduction copy of variable
in outer ctx if if exists. Check if non-existent in field_map before
installing OMP_CLAUSE_IN_REDUCTION decl.
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.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index a64b7f5aa10..8179b5aa8bc 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1138,7 +1138,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;
@@ -1285,6 +1285,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;
}
@@ -1353,7 +1366,7 @@ 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 openacc = false, bool openmp_target = false)
{
bool error = false;
gfc_omp_clauses *c = gfc_get_omp_clauses ();
@@ -2057,8 +2070,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,
@@ -3496,7 +3509,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,
+ (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 e55e0c81868..08483951066 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -6391,12 +6391,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 27a513e2539..8c0141b5cae 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1317,9 +1317,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (is_omp_target (ctx->stmt))
{
tree at = decl;
+ omp_context *scan_ctx = ctx;
if (ctx->outer)
- scan_omp_op (&at, ctx->outer);
- tree nt = omp_copy_decl_1 (at, ctx);
+ {
+ scan_omp_op (&at, ctx->outer);
+ scan_ctx = ctx->outer;
+ }
+ tree nt = omp_copy_decl_1 (at, scan_ctx);
splay_tree_insert (ctx->field_map,
(splay_tree_key) &DECL_CONTEXT (decl),
(splay_tree_value) nt);
@@ -1339,7 +1343,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (!is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)))
{
by_ref = use_pointer_for_field (decl, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION
+ && !splay_tree_lookup (ctx->field_map,
+ (splay_tree_key) decl))
install_var_field (decl, by_ref, 3, ctx);
}
install_var_local (decl, ctx);
diff --git a/gcc/testsuite/gfortran.dg/gomp/reduction4.f90
b/gcc/testsuite/gfortran.dg/gomp/reduction4.f90
index 52d504bac71..71b4231f315 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 00000000000..68512e223ac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-in-reduction-1.f90
@@ -0,0 +1,33 @@
+! { dg-do run }
+
+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
+
+program main
+ integer :: x, y
+
+ x = 1
+ y = 1
+
+ call foo (x, y)
+
+ if (x .ne. 11) stop 1
+ if (y .ne. 21) stop 2
+
+end program main
