* [PATCH, OpenMP, Fortran] Support in_reduction for Fortran
@ 2021-09-17 11:57 Chung-Lin Tang
2021-09-17 16:11 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Chung-Lin Tang @ 2021-09-17 11:57 UTC (permalink / raw)
To: gcc-patches, Fortran List, Jakub Jelinek, Tobias Burnus, Catherine Moore
[-- Attachment #1: Type: text/plain, Size: 1576 bytes --]
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 <cltang@codesourcery.com>
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.
[-- Attachment #2: fortran-in_reduction.patch --]
[-- Type: text/plain, Size: 6029 bytes --]
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
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH, OpenMP, Fortran] Support in_reduction for Fortran
2021-09-17 11:57 [PATCH, OpenMP, Fortran] Support in_reduction for Fortran Chung-Lin Tang
@ 2021-09-17 16:11 ` Jakub Jelinek
2021-10-19 13:03 ` [PATCH, v2, " Chung-Lin Tang
0 siblings, 1 reply; 4+ messages in thread
From: Jakub Jelinek @ 2021-09-17 16:11 UTC (permalink / raw)
To: Chung-Lin Tang; +Cc: gcc-patches, Fortran List, Tobias Burnus, Catherine Moore
On Fri, Sep 17, 2021 at 07:57:38PM +0800, Chung-Lin Tang wrote:
> 2021-09-17 Chung-Lin Tang <cltang@codesourcery.com>
>
> 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.
> @@ -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)
The ()s around op == EXEC_OMP_TARGET are unnecessary.
> --- 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);
If this was missing, then I'm afraid we lack a lot of testsuite coverage for
Fortran task reductions. It doesn't need to be covered in this patch, but would be
good to cover it incrementally. Because the above means nothing with
taskgroup with task_reduction clause(s) could work properly at runtime.
> --- 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);
You're right that the var remembered with &DECL_CONTEXT (whatever) key is
used outside of the target construct rather than inside of it.
So, if ctx->outer is non-NULL, it seems right to create the var in that
outer context. But, if ctx->outer is NULL, which can happen if the
target construct is orphaned, consider e.g.
extern int &x;
extern int &y;
void
foo ()
{
#pragma omp target in_reduction (+: x, y)
{
x = x + 8;
y = y + 16;
}
}
void
bar ()
{
#pragma omp taskgroup task_reduction (+: x, y)
foo ();
}
then those artificial decls (copies of x and y) should appear
to be at the function scope and not inside of the target region.
Therefore, I wonder if omp_copy_decl_2 shouldn't do the
DECL_CONTEXT (copy) = current_function_decl;
DECL_CHAIN (copy) = ctx->block_vars;
ctx->block_vars = copy;
(the last one can be moved next to the others) only if ctx != NULL
and otherwise call gimple_add_tmp_var (copy); instead
and then just call omp_copy_decl_1 at that spot with unconditional
ctx->outer.
Also, this isn't the only place that should have such a change,
there is also
if (ctx->outer)
scan_omp_op (&at, ctx->outer);
tree nt = omp_copy_decl_1 (at, ctx);
splay_tree_insert (ctx->field_map,
(splay_tree_key) &DECL_CONTEXT (t),
(splay_tree_value) nt);
a few lines above this and I'd expect that it should be (at, ctx->outer)
as well.
> @@ -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);
When exactly do you need this? It doesn't trigger on the new libgomp
testcase...
> --- /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
Again, something that can be dealt incrementally, but the
testsuite coverage of
https://gcc.gnu.org/pipermail/gcc-patches/2021-June/573600.html
was larger than this. Would be nice e.g. to cover both scalar vars
and array sections/arrays, parameters passed by reference as in the
above testcase, but also something that isn't a reference (either a local
variable or dummy parameter with VALUE, etc.
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
* [PATCH, v2, OpenMP, Fortran] Support in_reduction for Fortran
2021-09-17 16:11 ` Jakub Jelinek
@ 2021-10-19 13:03 ` Chung-Lin Tang
2021-10-19 14:05 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Chung-Lin Tang @ 2021-10-19 13:03 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches, Fortran List, Tobias Burnus, Catherine Moore
[-- Attachment #1: Type: text/plain, Size: 6846 bytes --]
Hi Jakub,
On 2021/9/18 12:11 AM, Jakub Jelinek wrote:
>> @@ -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)
>
> The ()s around op == EXEC_OMP_TARGET are unnecessary.
Fixed.
>> --- 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);
>
> If this was missing, then I'm afraid we lack a lot of testsuite coverage for
> Fortran task reductions. It doesn't need to be covered in this patch, but would be
> good to cover it incrementally. Because the above means nothing with
> taskgroup with task_reduction clause(s) could work properly at runtime.
Actually, the testcases do somewhat exercise taskgroup task_reductions, but like you
said, only lightly.
>> --- 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);
>
> You're right that the var remembered with &DECL_CONTEXT (whatever) key is
> used outside of the target construct rather than inside of it.
> So, if ctx->outer is non-NULL, it seems right to create the var in that
> outer context. But, if ctx->outer is NULL, which can happen if the
> target construct is orphaned, consider e.g.
> extern int &x;
> extern int &y;
>
> void
> foo ()
> {
> #pragma omp target in_reduction (+: x, y)
> {
> x = x + 8;
> y = y + 16;
> }
> }
>
> void
> bar ()
> {
> #pragma omp taskgroup task_reduction (+: x, y)
> foo ();
> }
> then those artificial decls (copies of x and y) should appear
> to be at the function scope and not inside of the target region.
>
> Therefore, I wonder if omp_copy_decl_2 shouldn't do the
> DECL_CONTEXT (copy) = current_function_decl;
> DECL_CHAIN (copy) = ctx->block_vars;
> ctx->block_vars = copy;
> (the last one can be moved next to the others) only if ctx != NULL
> and otherwise call gimple_add_tmp_var (copy); instead
> and then just call omp_copy_decl_1 at that spot with unconditional
> ctx->outer.
I see what you mean. I tried gimple_add_tmp_var but didn't work due to a
!DECL_SEEN_IN_BIND_EXPR_P() assert fail, but record_vars() appears to work.
> Also, this isn't the only place that should have such a change,
> there is also
> if (ctx->outer)
> scan_omp_op (&at, ctx->outer);
> tree nt = omp_copy_decl_1 (at, ctx);
> splay_tree_insert (ctx->field_map,
> (splay_tree_key) &DECL_CONTEXT (t),
> (splay_tree_value) nt);
> a few lines above this and I'd expect that it should be (at, ctx->outer)
> as well.
Fixed.
>> @@ -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);
>
> When exactly do you need this? It doesn't trigger on the new libgomp
> testcase...
I remember there was a testcase with triggered an ICE without this, but for some
reason can't find it anymore. I don't have any more evidence this is needed, so
removed now.
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.fortran/target-in-reduction-1.f90
>> @@ -0,0 +1,33 @@
>> +! { dg-do run }
>> +
>> +subroutine foo (x, y)
...
>> + if (x .ne. 11) stop 1
>> + if (y .ne. 21) stop 2
>> +
>> +end program main
>
> Again, something that can be dealt incrementally, but the
> testsuite coverage of
> https://gcc.gnu.org/pipermail/gcc-patches/2021-June/573600.html
> was larger than this. Would be nice e.g. to cover both scalar vars
> and array sections/arrays, parameters passed by reference as in the
> above testcase, but also something that isn't a reference (either a local
> variable or dummy parameter with VALUE, etc.
>
> Jakub
I have expanded target-in-reduction-1.f90 to cover local variables and
VALUE passed parameters. Array sections in reductions appear to be still
not supported by the Fortran FE in general (Tobias plans to work on that later).
I also added another target-in-reduction-2.f90 testcase that tests the "orphaned"
case in Fortran, where the task/target-in_reduction is in another separate subroutine.
Tested without regressions on trunk, is this okay to commit?
Thanks,
Chung-Lin
2021-10-19 Chung-Lin Tang <cltang@codesourcery.com>
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.
[-- Attachment #2: fortran-in_reduction-v2.patch --]
[-- Type: text/plain, Size: 8181 bytes --]
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 6a4ca2868f8..210fb06dbec 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,
@@ -3512,7 +3525,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 d234d1b070f..56efe195257 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -6405,12 +6405,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 057b7ae4866..15e4424b0bc 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 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..f9acb711e67
--- /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 00000000000..7f2e16b534b
--- /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
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH, v2, OpenMP, Fortran] Support in_reduction for Fortran
2021-10-19 13:03 ` [PATCH, v2, " Chung-Lin Tang
@ 2021-10-19 14:05 ` Jakub Jelinek
0 siblings, 0 replies; 4+ messages in thread
From: Jakub Jelinek @ 2021-10-19 14:05 UTC (permalink / raw)
To: Chung-Lin Tang; +Cc: Tobias Burnus, gcc-patches, Fortran List
On Tue, Oct 19, 2021 at 09:03:06PM +0800, Chung-Lin Tang wrote:
> 2021-10-19 Chung-Lin Tang <cltang@codesourcery.com>
>
> 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.
LGTM, thanks.
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2021-10-19 14:05 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-09-17 11:57 [PATCH, OpenMP, Fortran] Support in_reduction for Fortran Chung-Lin Tang
2021-09-17 16:11 ` Jakub Jelinek
2021-10-19 13:03 ` [PATCH, v2, " Chung-Lin Tang
2021-10-19 14:05 ` Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).