From: Cesar Philippidis Date: Fri, 29 Apr 2016 17:42:04 +0000 (-0700) Subject: re PR middle-end/70626 (bogus results in 'acc parallel loop' reductions) X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=e7ff0319f3736617c70742d8233d73faad523aa3;p=gcc.git re PR middle-end/70626 (bogus results in 'acc parallel loop' reductions) gcc/c-family/ PR middle-end/70626 * c-common.h (c_oacc_split_loop_clauses): Add boolean argument. * c-omp.c (c_oacc_split_loop_clauses): Use it to duplicate reduction clauses in acc parallel loops. gcc/c/ PR middle-end/70626 * c-parser.c (c_parser_oacc_loop): Don't augment mask with OACC_LOOP_CLAUSE_MASK. (c_parser_oacc_kernels_parallel): Update call to c_oacc_split_loop_clauses. gcc/cp/ PR middle-end/70626 * parser.c (cp_parser_oacc_loop): Don't augment mask with OACC_LOOP_CLAUSE_MASK. (cp_parser_oacc_kernels_parallel): Update call to c_oacc_split_loop_clauses. gcc/fortran/ PR middle-end/70626 * trans-openmp.c (gfc_trans_oacc_combined_directive): Duplicate the reduction clause in both parallel and loop directives. gcc/testsuite/ PR middle-end/70626 * c-c++-common/goacc/combined-reduction.c: New test. * gfortran.dg/goacc/reduction-2.f95: Add check for kernels reductions. libgomp/ PR middle-end/70626 * testsuite/libgomp.oacc-c++/template-reduction.C: Adjust test. * testsuite/libgomp.oacc-c-c++-common/combined-reduction.c: New test. * testsuite/libgomp.oacc-fortran/combined-reduction.f90: New test. From-SVN: r235651 --- diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index 2a740a57e7b..834950ef422 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,10 @@ +2016-04-29 Cesar Philippidis + + PR middle-end/70626 + * c-common.h (c_oacc_split_loop_clauses): Add boolean argument. + * c-omp.c (c_oacc_split_loop_clauses): Use it to duplicate + reduction clauses in acc parallel loops. + 2016-04-29 Marek Polacek PR c/70852 diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 3a7805f10a5..99fa3ab2937 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1278,7 +1278,7 @@ extern bool c_omp_check_loop_iv (tree, tree, walk_tree_lh); extern bool c_omp_check_loop_iv_exprs (location_t, tree, tree, tree, tree, walk_tree_lh); extern tree c_finish_oacc_wait (location_t, tree, tree); -extern tree c_oacc_split_loop_clauses (tree, tree *); +extern tree c_oacc_split_loop_clauses (tree, tree *, bool); extern void c_omp_split_clauses (location_t, enum tree_code, omp_clause_mask, tree, tree *); extern tree c_omp_declare_simd_clauses_to_numbers (tree, tree); diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index 5469d0d5625..be401bbb6b4 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -861,9 +861,10 @@ c_omp_check_loop_iv_exprs (location_t stmt_loc, tree declv, tree decl, #pragma acc parallel loop */ tree -c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) +c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses, + bool is_parallel) { - tree next, loop_clauses; + tree next, loop_clauses, nc; loop_clauses = *not_loop_clauses = NULL_TREE; for (; clauses ; clauses = next) @@ -882,7 +883,23 @@ c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_PRIVATE: + OMP_CLAUSE_CHAIN (clauses) = loop_clauses; + loop_clauses = clauses; + break; + + /* Reductions must be duplicated on both constructs. */ case OMP_CLAUSE_REDUCTION: + if (is_parallel) + { + nc = build_omp_clause (OMP_CLAUSE_LOCATION (clauses), + OMP_CLAUSE_REDUCTION); + OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_DECL (clauses); + OMP_CLAUSE_REDUCTION_CODE (nc) + = OMP_CLAUSE_REDUCTION_CODE (clauses); + OMP_CLAUSE_CHAIN (nc) = *not_loop_clauses; + *not_loop_clauses = nc; + } + OMP_CLAUSE_CHAIN (clauses) = loop_clauses; loop_clauses = clauses; break; diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 5161f7d56df..f02fbdb3399 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,11 @@ +2016-04-29 Cesar Philippidis + + PR middle-end/70626 + * c-parser.c (c_parser_oacc_loop): Don't augment mask with + OACC_LOOP_CLAUSE_MASK. + (c_parser_oacc_kernels_parallel): Update call to + c_oacc_split_loop_clauses. + 2016-04-28 Andrew MacLeod * c-array-notation.c (fix_builtin_array_notation_fn): Fix final diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 36c44ab5b56..832b8dda486 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -13828,6 +13828,8 @@ static tree c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, omp_clause_mask mask, tree *cclauses, bool *if_p) { + bool is_parallel = ((mask >> PRAGMA_OACC_CLAUSE_REDUCTION) & 1) == 1; + strcat (p_name, " loop"); mask |= OACC_LOOP_CLAUSE_MASK; @@ -13835,7 +13837,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, cclauses == NULL); if (cclauses) { - clauses = c_oacc_split_loop_clauses (clauses, cclauses); + clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel); if (*cclauses) *cclauses = c_finish_omp_clauses (*cclauses, false); if (clauses) @@ -13930,8 +13932,6 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, if (strcmp (p, "loop") == 0) { c_parser_consume_token (parser); - mask |= OACC_LOOP_CLAUSE_MASK; - tree block = c_begin_omp_parallel (); tree clauses; c_parser_oacc_loop (loc, parser, p_name, mask, &clauses, if_p); diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index d8f35af2e30..bb2b7cf5d12 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,11 @@ +2016-04-29 Cesar Philippidis + + PR middle-end/70626 + * parser.c (cp_parser_oacc_loop): Don't augment mask with + OACC_LOOP_CLAUSE_MASK. + (cp_parser_oacc_kernels_parallel): Update call to + c_oacc_split_loop_clauses. + 2016-04-28 Jason Merrill Implement C++17 [[nodiscard]] attribute. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7fb3c014ace..ded0dee6b5f 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -35408,6 +35408,8 @@ static tree cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, omp_clause_mask mask, tree *cclauses, bool *if_p) { + bool is_parallel = ((mask >> PRAGMA_OACC_CLAUSE_REDUCTION) & 1) == 1; + strcat (p_name, " loop"); mask |= OACC_LOOP_CLAUSE_MASK; @@ -35415,7 +35417,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, cclauses == NULL); if (cclauses) { - clauses = c_oacc_split_loop_clauses (clauses, cclauses); + clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel); if (*cclauses) *cclauses = finish_omp_clauses (*cclauses, false); if (clauses) @@ -35508,8 +35510,6 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, if (strcmp (p, "loop") == 0) { cp_lexer_consume_token (parser->lexer); - mask |= OACC_LOOP_CLAUSE_MASK; - tree block = begin_omp_parallel (); tree clauses; cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, &clauses, diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index 086a516b6bc..6a6b85ea1fc 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,9 @@ +2016-04-29 Cesar Philippidis + + PR middle-end/70626 + * trans-openmp.c (gfc_trans_oacc_combined_directive): Duplicate + the reduction clause in both parallel and loop directives. + 2016-04-18 Michael Matz * trans-io.c (gfc_build_io_library_fndecls): Use SET_TYPE_ALIGN. diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index a905ca607ae..c2d89eb3eff 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -3497,7 +3497,8 @@ gfc_trans_oacc_combined_directive (gfc_code *code) construct_clauses.independent = false; construct_clauses.tile_list = NULL; construct_clauses.lists[OMP_LIST_PRIVATE] = NULL; - construct_clauses.lists[OMP_LIST_REDUCTION] = NULL; + if (construct_code == OACC_KERNELS) + construct_clauses.lists[OMP_LIST_REDUCTION] = NULL; oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses, code->loc); } diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index f8634553277..addef50175a 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2016-04-29 Cesar Philippidis + + PR middle-end/70626 + * c-c++-common/goacc/combined-reduction.c: New test. + * gfortran.dg/goacc/reduction-2.f95: Add check for kernels reductions. + 2016-04-29 H.J. Lu * gcc.target/i386/pr70155-1.c: Check for nonexistence of the diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c new file mode 100644 index 00000000000..ecf23f59d66 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c @@ -0,0 +1,29 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenacc -fdump-tree-gimple" } */ + +#include + +int +main () +{ + int i, v1 = 0, n = 100; + +#pragma acc parallel loop reduction(+:v1) + for (i = 0; i < n; i++) + v1++; + + assert (v1 == n); + +#pragma acc kernels loop reduction(+:v1) + for (i = 0; i < n; i++) + v1++; + + assert (v1 == n); + + return 0; +} + +/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 b/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 index 929fb0e81e9..4d40998958c 100644 --- a/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 @@ -15,7 +15,7 @@ subroutine foo () !$acc end kernels loop end subroutine -! { dg-final { scan-tree-dump-times "target oacc_parallel firstprivate.a." 1 "gimple" } } +! { dg-final { scan-tree-dump-times "target oacc_parallel reduction..:a. map.tofrom.a." 1 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop private.p. reduction..:a." 1 "gimple" } } ! { dg-final { scan-tree-dump-times "target oacc_kernels map.force_tofrom:a .len: 4.." 1 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop private.k. reduction..:a." 1 "gimple" } } diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 9dfc7f37f0d..351c2392530 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,10 @@ +2016-04-29 Cesar Philippidis + + PR middle-end/70626 + * testsuite/libgomp.oacc-c++/template-reduction.C: Adjust test. + * testsuite/libgomp.oacc-c-c++-common/combined-reduction.c: New test. + * testsuite/libgomp.oacc-fortran/combined-reduction.f90: New test. + 2016-04-21 Alexander Monakov * plugin/plugin-nvptx.c (map_fini): Make cuMemFreeHost error diff --git a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C b/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C index fb5924c9b51..6c85fba7a92 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C +++ b/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C @@ -7,7 +7,7 @@ sum (T array[]) { T s = 0; -#pragma acc parallel loop num_gangs (10) gang reduction (+:s) copy (s, array[0:n]) +#pragma acc parallel loop num_gangs (10) gang reduction (+:s) copy (array[0:n]) for (int i = 0; i < n; i++) s += array[i]; @@ -25,7 +25,7 @@ sum () for (int i = 0; i < n; i++) array[i] = i+1; -#pragma acc parallel loop num_gangs (10) gang reduction (+:s) copy (s) +#pragma acc parallel loop num_gangs (10) gang reduction (+:s) for (int i = 0; i < n; i++) s += array[i]; @@ -43,7 +43,7 @@ async_sum (T array[]) for (int i = 0; i < n; i++) array[i] = i+1; -#pragma acc parallel loop num_gangs (10) gang reduction (+:s) present (array[0:n]) copy (s) async wait (1) +#pragma acc parallel loop num_gangs (10) gang reduction (+:s) present (array[0:n]) async wait (1) for (int i = 0; i < n; i++) s += array[i]; @@ -59,7 +59,7 @@ async_sum (int c) { T s = 0; -#pragma acc parallel loop num_gangs (10) gang reduction (+:s) copy(s) firstprivate (c) async wait (1) +#pragma acc parallel loop num_gangs (10) gang reduction (+:s) firstprivate (c) async wait (1) for (int i = 0; i < n; i++) s += i+c; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-reduction.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-reduction.c new file mode 100644 index 00000000000..b5ce4ed7b3d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-reduction.c @@ -0,0 +1,23 @@ +/* Test a combined acc parallel loop reduction. */ + +/* { dg-do run } */ + +#include + +int +main () +{ + int i, v1 = 0, v2 = 0, n = 100; + +#pragma acc parallel loop reduction(+:v1, v2) + for (i = 0; i < n; i++) + { + v1++; + v2++; + } + + assert (v1 == n); + assert (v2 == n); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/combined-reduction.f90 b/libgomp/testsuite/libgomp.oacc-fortran/combined-reduction.f90 new file mode 100644 index 00000000000..d3a61b57f13 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/combined-reduction.f90 @@ -0,0 +1,19 @@ +! Test a combined acc parallel loop reduction. + +! { dg-do run } + +program test + implicit none + integer i, n, var + + n = 100 + var = 0 + + !$acc parallel loop reduction(+:var) + do i = 1, 100 + var = var + 1 + end do + !$acc end parallel loop + + if (var .ne. n) call abort +end program test