This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

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


Hi!

On behalf of Gergő (who doesn't have write access yet) I've pushed the
attached to openacc-gcc-8-branch.


Grüße
 Thomas


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".

In gcc/testsuite/c-c++-common/goacc/reduction-6.c, we remove the erroneous
reductions on variable b; adding a reduction clause to make it compile cleanly
would make it a duplicate of the test for variable c.

	gcc/
	* omp-low.c (struct omp_context): New fields
	local_reduction_clauses, outer_reduction_clauses.
	(new_omp_context): Initialize these.
	(scan_sharing_clauses): Record reduction clauses on OpenACC
	constructs.
	(scan_omp_for): Check reduction clauses for incorrect nesting.
	gcc/testsuite/
	* 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.
	libgomp/
	* 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.
---
 gcc/ChangeLog.openacc                         |   9 +
 gcc/omp-low.c                                 | 105 ++++
 gcc/testsuite/ChangeLog.openacc               |   7 +
 .../goacc/nested-reductions-fail.c            | 492 ++++++++++++++++++
 .../c-c++-common/goacc/nested-reductions.c    | 420 +++++++++++++++
 .../c-c++-common/goacc/reduction-6.c          |  11 -
 libgomp/ChangeLog.openacc                     |  12 +
 .../par-loop-comb-reduction-1.c               |   2 +-
 .../par-loop-comb-reduction-2.c               |   2 +-
 .../par-loop-comb-reduction-3.c               |   2 +-
 .../par-loop-comb-reduction-4.c               |   2 +-
 11 files changed, 1049 insertions(+), 15 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions.c

diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc
index e57971ce4cf2..597362505bb3 100644
--- a/gcc/ChangeLog.openacc
+++ b/gcc/ChangeLog.openacc
@@ -1,3 +1,12 @@
+2018-12-20  Gergö Barany  <gergo@codesourcery.com>
+
+	* omp-low.c (struct omp_context): New fields
+	local_reduction_clauses, outer_reduction_clauses.
+	(new_omp_context): Initialize these.
+	(scan_sharing_clauses): Record reduction clauses on OpenACC
+	constructs.
+	(scan_omp_for): Check reduction clauses for incorrect nesting.
+
 2018-12-20  Julian Brown  <julian@codesourcery.com>
 	    Maciej W. Rozycki  <macro@codesourcery.com>
 
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 967916521001..6b7b23e6909b 100644
--- 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;
     }
   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)
-- 
2.17.1


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]