[og8] Report errors on missing OpenACC reduction clauses in nested reductions

Thomas Schwinge thomas@codesourcery.com
Mon Apr 20 14:53:01 GMT 2020


Hi Frederik!

As you've been the last one to work on this code and get it into GCC
master branch (commit 5d183d1740d8d8b84991f186ce4d992ee799536f "Warn
about inconsistent OpenACC nested reduction clauses"), you get to look
into the following issue.

..., and yes, I'm aware that you have a follow-up patch which is
reworking this code, and which still needs to be reviewed --
<http://mid.mail-archive.com/7ff53a86-d2ec-21ea-d0f4-4e0b13986771@codesourcery.com>
-- but let's please first have a look at the current code in GCC master
branch, and address what might be a bug.

On 2018-12-20T15:28:33+0100, I wrote:
> On behalf of Gergő (who doesn't have write access yet) I've pushed the
> attached to openacc-gcc-8-branch.

> From a9e48066198ffb1e7bc2b137167a61a6cb47748c Mon Sep 17 00:00:00 2001
> From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@codesourcery.com>
> Date: Thu, 20 Dec 2018 15:07:34 +0100
> Subject: [PATCH] Report errors on missing OpenACC reduction clauses in nested
>  reductions
>
> ..., as suggested by OpenACC 2.6, 2.9.11. "reduction clause".

> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -129,6 +129,12 @@ struct omp_context
>
>    /* Hash map of dynamic arrays in this context.  */
>    hash_map<tree_operand_hash, tree> *dynamic_arrays;
> +
> +  /* A tree_list of the reduction clauses in this context.  */
> +  tree local_reduction_clauses;
> +
> +  /* A tree_list of the reduction clauses in outer contexts.  */
> +  tree outer_reduction_clauses;
>  };
>
>  static splay_tree all_contexts;
> @@ -1040,6 +1046,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
>        ctx->cb = outer_ctx->cb;
>        ctx->cb.block = NULL;
>        ctx->depth = outer_ctx->depth + 1;
> +      ctx->local_reduction_clauses = NULL;
> +      ctx->outer_reduction_clauses = ctx->outer_reduction_clauses;

Via <https://gcc.gnu.org/PR94629> "10 issues located by the PVS-studio
static analyzer" (so please reference that one on any patch submission),
on <https://habr.com/en/company/pvs-studio/blog/497640/> in "Fragment N3,
Assigning a variable to itself", we find this latter assignment qualified
as "very strange to assign a variable to itself".

Probably that should've been 'outer_ctx' instead of 'ctx'?  But: why/how
then does the current algorith still work despite this error?

(I haven't looked at the code very much.)


Grüße
 Thomas


>      }
>    else
>      {
> @@ -1053,6 +1061,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
>        ctx->cb.eh_lp_nr = 0;
>        ctx->cb.transform_call_graph_edges = CB_CGE_MOVE;
>        ctx->depth = 1;
> +      ctx->local_reduction_clauses = NULL;
> +      ctx->outer_reduction_clauses = NULL;
>      }
>
>    ctx->cb.decl_map = new hash_map<tree, tree>;
> @@ -1276,6 +1286,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
>         goto do_private;
>
>       case OMP_CLAUSE_REDUCTION:
> +          if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx))
> +            ctx->local_reduction_clauses
> +           = tree_cons (NULL, c, ctx->local_reduction_clauses);
>         decl = OMP_CLAUSE_DECL (c);
>         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
>             && TREE_CODE (decl) == MEM_REF)
> @@ -2458,6 +2471,98 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
>         gimple_omp_for_set_clauses (stmt, clauses);
>         check_oacc_kernel_gwv (stmt, ctx);
>       }
> +
> +      /* Collect all variables named in reductions on this loop.  Ensure
> +         that, if this loop has a reduction on some variable v, and there is
> +         a reduction on v somewhere in an outer context, then there is a
> +         reduction on v on all intervening loops as well.  */
> +      tree local_reduction_clauses = NULL;
> +      for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
> +        {
> +          if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
> +            local_reduction_clauses
> +           = tree_cons (NULL, c, local_reduction_clauses);
> +        }
> +      if (ctx->outer_reduction_clauses == NULL && ctx->outer != NULL)
> +        ctx->outer_reduction_clauses
> +       = chainon (unshare_expr (ctx->outer->local_reduction_clauses),
> +                  ctx->outer->outer_reduction_clauses);
> +      tree outer_reduction_clauses = ctx->outer_reduction_clauses;
> +      tree local_iter = local_reduction_clauses;
> +      for (; local_iter; local_iter = TREE_CHAIN (local_iter))
> +        {
> +          tree local_clause = TREE_VALUE (local_iter);
> +          tree local_var = OMP_CLAUSE_DECL (local_clause);
> +          tree_code local_op = OMP_CLAUSE_REDUCTION_CODE (local_clause);
> +          bool have_outer_reduction = false;
> +          tree ctx_iter = outer_reduction_clauses;
> +          for (; ctx_iter; ctx_iter = TREE_CHAIN (ctx_iter))
> +            {
> +              tree outer_clause = TREE_VALUE (ctx_iter);
> +              tree outer_var = OMP_CLAUSE_DECL (outer_clause);
> +              tree_code outer_op = OMP_CLAUSE_REDUCTION_CODE (outer_clause);
> +              if (outer_var == local_var && outer_op != local_op)
> +                {
> +                  error_at (gimple_location (stmt),
> +                            "conflicting reduction operations for %qE",
> +                            local_var);
> +                  inform (OMP_CLAUSE_LOCATION (outer_clause),
> +                          "location of the previous reduction for %qE",
> +                          outer_var);
> +                  /* Change this operation to be equal to the outer one.
> +                     This is meant to suppress spurious errors; for example,
> +                     in nested +, -, + reductions, we would generate errors
> +                     for both the change from + to - and from - to +.  */
> +                  OMP_CLAUSE_REDUCTION_CODE (local_clause) = outer_op;
> +                  /* Also change the location so that in nested +, -, -
> +                     reductions, the second error message also refers to the
> +                     outermost + reduction.  */
> +                  OMP_CLAUSE_LOCATION (local_clause)
> +                 = OMP_CLAUSE_LOCATION (outer_clause);
> +                }
> +              if (outer_var == local_var)
> +                {
> +                  have_outer_reduction = true;
> +                  break;
> +                }
> +            }
> +          if (have_outer_reduction)
> +            {
> +              /* There is a reduction on outer_var both on this loop and on
> +                 some enclosing loop.  Walk up the context tree until such a
> +                 loop with a reduction on outer_var is found, and complain
> +                 about all intervening loops that do not have such a
> +                 reduction.  */
> +              struct omp_context *curr_loop = ctx->outer;
> +              bool found = false;
> +              while (curr_loop != NULL)
> +                {
> +                  tree curr_iter = curr_loop->local_reduction_clauses;
> +                  for (; curr_iter; curr_iter = TREE_CHAIN (curr_iter))
> +                    {
> +                      tree curr_clause = TREE_VALUE (curr_iter);
> +                      tree curr_var = OMP_CLAUSE_DECL (curr_clause);
> +                      if (curr_var == local_var)
> +                        {
> +                          found = true;
> +                          break;
> +                        }
> +                    }
> +                  if (!found)
> +                    error_at (gimple_location (curr_loop->stmt),
> +                              "nested loop in reduction needs "
> +                              "reduction clause for %qE",
> +                              local_var);
> +                  else
> +                    break;
> +                  curr_loop = curr_loop->outer;
> +                }
> +            }
> +        }
> +      ctx->local_reduction_clauses = local_reduction_clauses;
> +      ctx->outer_reduction_clauses
> +     = chainon (unshare_expr (ctx->local_reduction_clauses),
> +                ctx->outer_reduction_clauses);
>      }
>
>    scan_sharing_clauses (clauses, ctx);
> diff --git a/gcc/testsuite/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc
> index b1758eca0b21..4af31e5e1060 100644
> --- a/gcc/testsuite/ChangeLog.openacc
> +++ b/gcc/testsuite/ChangeLog.openacc
> @@ -1,3 +1,10 @@
> +2018-12-20  Gergö Barany  <gergo@codesourcery.com>
> +         Thomas Schwinge  <thomas@codesourcery.com>
> +
> +     * c-c++-common/goacc/nested-reductions-fail.c: New test.
> +     * c-c++-common/goacc/nested-reductions.c: New test.
> +     * c-c++-common/goacc/reduction-6.c: Adjust.
> +
>  2018-12-20  Maciej W. Rozycki  <macro@codesourcery.com>
>
>       * c-c++-common/goacc/serial-dims.c: New test.
> diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
> new file mode 100644
> index 000000000000..a642dd038709
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
> @@ -0,0 +1,492 @@
> +/* Test erroneous cases of nested reduction loops.  */
> +
> +void acc_parallel (void)
> +{
> +  int i, j, k, l, sum, diff;
> +
> +  #pragma acc parallel
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but using a combined parallel loop construct.  */
> +
> +void acc_parallel_loop (void)
> +{
> +  int i, j, k, l, sum, diff;
> +
> +  #pragma acc parallel loop
> +  for (int h = 0; h < 10; ++h)
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but now the outermost reduction clause is on
> +   the parallel region, not the outermost loop.  */
> +void acc_parallel_reduction (void)
> +{
> +  int i, j, k, l, sum, diff;
> +
> +  #pragma acc parallel reduction(+:sum)
> +  {
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but using a combined parallel loop construct, and
> +   the outermost reduction clause is on that one, not the outermost loop.  */
> +void acc_parallel_loop_reduction (void)
> +{
> +  int i, j, k, l, sum, diff;
> +
> +  #pragma acc parallel loop reduction(+:sum)
> +  for (int h = 0; h < 10; ++h)
> +  {
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but inside a routine construct.  */
> +#pragma acc routine gang
> +void acc_routine (void)
> +{
> +  int i, j, k, l, sum, diff;
> +
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +          #pragma acc loop reduction(+:sum)
> +          for (l = 0; l < 10; l++)
> +            sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +       #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
> +       for (l = 0; l < 10; l++)
> +         sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +void acc_kernels (void)
> +{
> +  int i, j, k, sum, diff;
> +
> +  /* FIXME:  No diagnostics are produced for these loops because reductions
> +     in kernels regions are not supported yet.  */
> +  #pragma acc kernels
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:diff)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(-:sum)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +  }
> +}
> diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions.c
> new file mode 100644
> index 000000000000..bff66527df58
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions.c
> @@ -0,0 +1,420 @@
> +/* Test cases of nested reduction loops that should compile cleanly.  */
> +
> +void acc_parallel (void)
> +{
> +  int i, j, k, sum, diff;
> +
> +  #pragma acc parallel
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop collapse(2) reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum)
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(-:diff)
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but using a combined parallel loop construct.  */
> +
> +void acc_parallel_loop (void)
> +{
> +  int i, j, k, l, sum, diff;
> +
> +  #pragma acc parallel loop
> +  for (int h = 0; h < 10; ++h)
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop collapse(2) reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but now the outermost reduction clause is on
> +   the parallel region, not the outermost loop.  */
> +
> +void acc_parallel_reduction (void)
> +{
> +  int i, j, k, sum, diff;
> +
> +  #pragma acc parallel reduction(+:sum)
> +  {
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum)
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(-:diff)
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum)
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(-:diff)
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum)
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but using a combined parallel loop construct, and
> +   the outermost reduction clause is on that one, not the outermost loop.  */
> +void acc_parallel_loop_reduction (void)
> +{
> +  int i, j, k, sum, diff;
> +
> +  #pragma acc parallel loop reduction(+:sum)
> +  for (int h = 0; h < 10; ++h)
> +  {
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop // { dg-warning "insufficient partitioning available to parallelize loop" }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +/* The same tests as above, but inside a routine construct.  */
> +#pragma acc routine gang
> +void acc_routine (void) // { dg-bogus "region is gang partitioned but does not contain gang partitioned code" "TODO" { xfail *-*-* } }
> +{
> +  int i, j, k, sum, diff;
> +
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop collapse(2) reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop collapse(2) reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum) // { dg-bogus "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum) reduction(-:diff)
> +    for (i = 0; i < 10; i++)
> +      {
> +        #pragma acc loop reduction(+:sum) // { dg-bogus "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(+:sum)
> +          for (k = 0; k < 10; k++)
> +            sum = 1;
> +
> +        #pragma acc loop reduction(-:diff) // { dg-bogus "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
> +        for (j = 0; j < 10; j++)
> +          #pragma acc loop reduction(-:diff)
> +          for (k = 0; k < 10; k++)
> +            diff = 1;
> +      }
> +  }
> +}
> +
> +void acc_kernels (void)
> +{
> +  int i, j, k, sum, diff;
> +
> +  /* FIXME:  These tests are not meaningful yet because reductions in
> +     kernels regions are not supported yet.  */
> +  #pragma acc kernels
> +  {
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +
> +    #pragma acc loop reduction(+:sum)
> +    for (i = 0; i < 10; i++)
> +      #pragma acc loop reduction(+:sum)
> +      for (j = 0; j < 10; j++)
> +        #pragma acc loop reduction(+:sum)
> +        for (k = 0; k < 10; k++)
> +          sum = 1;
> +  }
> +}
> diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-6.c b/gcc/testsuite/c-c++-common/goacc/reduction-6.c
> index 619f82b9d8b3..3c10b4dddaf2 100644
> --- a/gcc/testsuite/c-c++-common/goacc/reduction-6.c
> +++ b/gcc/testsuite/c-c++-common/goacc/reduction-6.c
> @@ -16,17 +16,6 @@ int foo (int N)
>        }
>    }
>
> -  #pragma acc parallel
> -  {
> -    #pragma acc loop reduction(+:b)
> -    for (int i = 0; i < N; i++)
> -      {
> -        #pragma acc loop
> -     for (int j = 0; j < N; j++)
> -       b += 1;
> -      }
> -  }
> -
>    #pragma acc parallel
>    {
>      #pragma acc loop reduction(+:c)
> diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc
> index b16008a77d19..b48453b3066a 100644
> --- a/libgomp/ChangeLog.openacc
> +++ b/libgomp/ChangeLog.openacc
> @@ -1,3 +1,15 @@
> +2018-12-20  Gergö Barany  <gergo@codesourcery.com>
> +         Thomas Schwinge  <thomas@codesourcery.com>
> +
> +     * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
> +     Add missing reduction clauses.
> +     * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
> +     Likewise.
> +     * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
> +     Likewise.
> +     * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
> +     Likewise.
> +
>  2018-12-20  Julian Brown  <julian@codesourcery.com>
>           Maciej W. Rozycki  <macro@codesourcery.com>
>
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
> index 5e82e1d350cc..91fe772045d9 100644
> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
> @@ -15,7 +15,7 @@ main (int argc, char *argv[])
>    #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
>      reduction(+:res) copy(res)
>    {
> -    #pragma acc loop gang
> +    #pragma acc loop gang reduction(+:res)
>      for (j = 0; j < 32; j++)
>        {
>       #pragma acc loop worker reduction(+:res)
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
> index a339f327956e..8a7bc7240708 100644
> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
> @@ -14,7 +14,7 @@ main (int argc, char *argv[])
>    #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
>                      reduction(^:res)
>    {
> -    #pragma acc loop gang
> +    #pragma acc loop gang reduction(^:res)
>      for (j = 0; j < 32; j++)
>        {
>       #pragma acc loop worker vector reduction(^:res)
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
> index 6369d7fbb330..eba5c6a544e8 100644
> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
> @@ -16,7 +16,7 @@ main (int argc, char *argv[])
>    #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
>      reduction(+:res) copy(res)
>    {
> -    #pragma acc loop gang
> +    #pragma acc loop gang reduction(+:res)
>      for (j = 0; j < 32; j++)
>        {
>       #pragma acc loop worker vector reduction(+:res)
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
> index 140c32263278..12b823f33ab1 100644
> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
> @@ -16,7 +16,7 @@ main (int argc, char *argv[])
>    #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
>      reduction(+:res) reduction(max:mres) copy(res, mres)
>    {
> -    #pragma acc loop gang
> +    #pragma acc loop gang reduction(+:res) reduction(max:mres)
>      for (j = 0; j < 32; j++)
>        {
>       #pragma acc loop worker vector reduction(+:res) reduction(max:mres)
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter


More information about the Gcc-patches mailing list