From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1709) id 923CD3858409; Wed, 20 Oct 2021 15:25:52 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 923CD3858409 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Chung-Lin Tang To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-4574] openmp: in_reduction support for Fortran X-Act-Checkin: gcc X-Git-Author: Chung-Lin Tang X-Git-Refname: refs/heads/master X-Git-Oldrev: 90454a900824d96e6d4eae557a809c9d986198d9 X-Git-Newrev: d98626bf451dea6a28a42d953f7d0bd7659ad4d5 Message-Id: <20211020152552.923CD3858409@sourceware.org> Date: Wed, 20 Oct 2021 15:25:52 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Wed, 20 Oct 2021 15:25:52 -0000 https://gcc.gnu.org/g:d98626bf451dea6a28a42d953f7d0bd7659ad4d5 commit r12-4574-gd98626bf451dea6a28a42d953f7d0bd7659ad4d5 Author: Chung-Lin Tang Date: Wed Oct 20 23:25:02 2021 +0800 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. Diff: --- gcc/fortran/openmp.c | 25 +++++-- gcc/fortran/trans-openmp.c | 9 ++- gcc/omp-low.c | 15 +++-- gcc/testsuite/gfortran.dg/gomp/reduction4.f90 | 2 +- .../libgomp.fortran/target-in-reduction-1.f90 | 78 ++++++++++++++++++++++ .../libgomp.fortran/target-in-reduction-2.f90 | 30 +++++++++ 6 files changed, 147 insertions(+), 12 deletions(-) diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 2a161f3304c..dcf22ac2c2f 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 4f1b3462959..aaeb950fb72 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 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