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]

[PATCH][og8] Update code and reduction tests for `serial' construct


This fixes a conflict between two recently committed patches to openacc-gcc-8-branch, Maciej's "Add OpenACC 2.6 `serial' construct support" and my "Report errors on missing OpenACC reduction clauses in nested reductions". The former renamed a function which caused the latter to no longer compile.

Additionally, new tests for OpenACC reductions in serial regions are added, and the existing ones separated out by region kind (parallel/kernels/serial).

OK for openacc-gcc-8-branch?


2018-12-21  Gergö Barany  <gergo@codesourcery.com>

	gcc/
	* omp-low.c (scan_sharing_clauses): Fix call to renamed function
	is_oacc_parallel.
	gcc/testsuite/c-c++-common/goacc/
	* nested-reductions-fail.c: Renamed to...
	* nested-reductions-parallel-fail.c: ...this file, with kernels tests...
	* nested-reductions-kernels-fail.c: ... moved to this new file.
	* nested-reductions-serial-fail.c: New test.
	* nested-reductions.c: Renamed to...
	* nested-reductions-parallel.c: ... this file, with kernels tests...
	* nested-reductions-kernels.c: ... moved to this new file.
	* nested-reductions-serial.c: New test.
>From 72098b852c0cee656f61395c04f9271a0a598761 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@codesourcery.com>
Date: Fri, 21 Dec 2018 00:08:09 -0800
Subject: [PATCH] [og8] Update code and reduction tests for `serial' construct

    gcc/
    * omp-low.c (scan_sharing_clauses): Fix call to renamed function
    is_oacc_parallel.
    gcc/testsuite/c-c++-common/goacc/
    * nested-reductions-fail.c: Renamed to...
    * nested-reductions-parallel-fail.c: ...this file, with kernels tests...
    * nested-reductions-kernels-fail.c: ... moved to this new file.
    * nested-reductions-serial-fail.c: New test.
    * nested-reductions.c: Renamed to...
    * nested-reductions-parallel.c: ... this file, with kernels tests...
    * nested-reductions-kernels.c: ... moved to this new file.
    * nested-reductions-serial.c: New test.
---
 gcc/ChangeLog.openacc                              |   5 +
 gcc/omp-low.c                                      |   2 +-
 gcc/testsuite/ChangeLog.openacc                    |  15 +
 .../c-c++-common/goacc/nested-reductions-fail.c    | 492 ---------------------
 .../goacc/nested-reductions-kernels-fail.c         | 273 ++++++++++++
 .../c-c++-common/goacc/nested-reductions-kernels.c | 227 ++++++++++
 .../goacc/nested-reductions-parallel-fail.c        | 447 +++++++++++++++++++
 .../goacc/nested-reductions-parallel.c             | 384 ++++++++++++++++
 .../goacc/nested-reductions-serial-fail.c          | 446 +++++++++++++++++++
 .../c-c++-common/goacc/nested-reductions-serial.c  | 391 ++++++++++++++++
 .../c-c++-common/goacc/nested-reductions.c         | 420 ------------------
 11 files changed, 2189 insertions(+), 913 deletions(-)
 delete mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-kernels-fail.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-kernels.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel-fail.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-serial-fail.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-serial.c
 delete mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions.c

diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc
index 5973625..718044c 100644
--- a/gcc/ChangeLog.openacc
+++ b/gcc/ChangeLog.openacc
@@ -1,3 +1,8 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+	* omp-low.c (scan_sharing_clauses): Fix call to renamed function
+	is_oacc_parallel.
+
 2018-12-20  Gergö Barany  <gergo@codesourcery.com>
 
 	* omp-low.c (struct omp_context): New fields
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 6b7b23e..72b6548 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1286,7 +1286,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	  goto do_private;
 
 	case OMP_CLAUSE_REDUCTION:
-          if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx))
+          if (is_gimple_omp_oacc (ctx->stmt))
             ctx->local_reduction_clauses
 	      = tree_cons (NULL, c, ctx->local_reduction_clauses);
 	  decl = OMP_CLAUSE_DECL (c);
diff --git a/gcc/testsuite/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc
index 4af31e5..473eb9d 100644
--- a/gcc/testsuite/ChangeLog.openacc
+++ b/gcc/testsuite/ChangeLog.openacc
@@ -1,3 +1,18 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+	* c-c++-common/goacc/nested-reductions-fail.c: Renamed to...
+	* c-c++-common/goacc/nested-reductions-parallel-fail.c: ...this file,
+	with kernels tests...
+	* c-c++-common/goacc/nested-reductions-kernels-fail.c: ... moved to this
+	new file.
+	* c-c++-common/goacc/nested-reductions-serial-fail.c: New test.
+	* c-c++-common/goacc/nested-reductions.c: Renamed to...
+	* c-c++-common/goacc/nested-reductions-parallel.c: ... this file, with
+	kernels tests...
+	* c-c++-common/goacc/nested-reductions-kernels.c: ... moved to this new
+	file.
+	* c-c++-common/goacc/nested-reductions-serial.c: New test.
+
 2018-12-20  Gergö Barany  <gergo@codesourcery.com>
 	    Thomas Schwinge  <thomas@codesourcery.com>
 
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
deleted file mode 100644
index a642dd0..0000000
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
+++ /dev/null
@@ -1,492 +0,0 @@
-/* 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-kernels-fail.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-kernels-fail.c
new file mode 100644
index 0000000..79545fb
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-kernels-fail.c
@@ -0,0 +1,273 @@
+/* Test erroneous cases of nested reduction loops in kernels regions,
+   corresponding to nested-reductions-parallel-fail.c.  */
+
+/* FIXME:  No diagnostics are produced for these loops because reductions
+   in kernels regions are not supported yet.  */
+
+void acc_kernels (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc kernels
+  {
+    #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 collapse(2)
+      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
+      for (j = 0; j < 10; j++)
+        #pragma acc loop
+        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)
+      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)
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum)
+      for (j = 0; j < 10; j++)
+        #pragma acc loop
+        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)
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+: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) reduction(-:diff)
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(-:diff)
+        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 (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 kernels loop construct.  */
+
+void acc_kernels_loop (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc kernels loop
+  for (int h = 0; h < 10; ++h)
+  {
+    #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 collapse(2)
+      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
+      for (j = 0; j < 10; j++)
+        #pragma acc loop
+        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)
+      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)
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum)
+      for (j = 0; j < 10; j++)
+        #pragma acc loop
+        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)
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+: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) reduction(-:diff)
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(-:diff)
+        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 (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 kernels loop construct, and
+   the outermost reduction clause is on that one, not the outermost loop.  */
+
+void acc_kernels_loop_reduction (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc kernels loop reduction(+:sum)
+  for (int h = 0; h < 10; ++h)
+  {
+    #pragma acc loop
+    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
+    for (i = 0; i < 10; i++)
+      #pragma acc loop collapse(2)
+      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
+    for (i = 0; i < 10; i++)
+      #pragma acc loop
+      for (j = 0; j < 10; j++)
+        #pragma acc loop
+        for (k = 0; k < 10; k++)
+          #pragma acc loop reduction(+:sum)
+          for (l = 0; l < 10; l++)
+            sum = 1;
+
+    #pragma acc loop
+    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 (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(max:sum)
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum)
+      for (j = 0; j < 10; j++)
+        #pragma acc loop
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum)
+	  for (l = 0; l < 10; l++)
+	    sum = 1;
+
+    #pragma acc loop reduction(max: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++)
+	  #pragma acc loop reduction(*:sum)
+	  for (l = 0; l < 10; l++)
+	    sum = 1;
+
+    #pragma acc loop reduction(-:diff)
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(-:diff)
+        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 (j = 0; j < 10; j++)
+          #pragma acc loop reduction(-:diff)
+          for (k = 0; k < 10; k++)
+            diff = 1;
+      }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-kernels.c
new file mode 100644
index 0000000..84863ff
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-kernels.c
@@ -0,0 +1,227 @@
+/* Test cases of nested reduction loops in kernels regions that should
+   compile cleanly, corresponding to nested-reductions-parallel.c.  */
+
+void acc_kernels (void)
+{
+  int i, j, k, sum, diff;
+
+  #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 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 kernels loop construct.  */
+
+void acc_kernels_loop (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc kernels 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)
+      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 kernels loop construct, and
+   the outermost reduction clause is on that one, not the outermost loop.  */
+
+void acc_kernels_loop_reduction (void)
+{
+  int i, j, k, sum, diff;
+
+  #pragma acc kernels 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)
+      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;
+      }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel-fail.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel-fail.c
new file mode 100644
index 0000000..aac7605
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel-fail.c
@@ -0,0 +1,447 @@
+/* Test erroneous cases of nested reduction loops in parallel regions.  See
+   also corresponding tests in nested-reductions-kernels-fail.c and
+   nested-reductions-serial-fail.c.  */
+
+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;
+      }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel.c
new file mode 100644
index 0000000..bfd1964
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-parallel.c
@@ -0,0 +1,384 @@
+/* Test cases of nested reduction loops in parallel regions and routines
+   that should compile cleanly.  See also corresponding tests in
+   nested-reductions-kernels.c and nested-reductions-serial.c.  */
+
+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;
+      }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-serial-fail.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-serial-fail.c
new file mode 100644
index 0000000..b526d09
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-serial-fail.c
@@ -0,0 +1,446 @@
+/* Test erroneous cases of nested reduction loops in serial regions,
+   corresponding to nested-reductions-parallel-fail.c.  */
+
+void acc_serial (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc serial
+  {
+    #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 serial loop construct.  */
+
+void acc_serial_loop (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc serial 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 serial region, not the outermost loop.  */
+void acc_serial_reduction (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc serial 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 serial loop construct, and
+   the outermost reduction clause is on that one, not the outermost loop.  */
+void acc_serial_loop_reduction (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc serial 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;
+      }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-serial.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-serial.c
new file mode 100644
index 0000000..05fea07
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-serial.c
@@ -0,0 +1,391 @@
+/* Test cases of nested reduction loops in serial regions that should
+   compile cleanly, corresponding to nested-reductions-parallel.c.  Since
+   serial constructs are not parallel, we must suppress some warnings about
+   insufficient parallelism.  */
+
+/* { dg-prune-output "insufficient partitioning available to parallelize loop" } */
+/* { dg-prune-output "region contains .* partitoned code but is not .* partitioned" } */
+
+
+void acc_serial (void)
+{
+  int i, j, k, sum, diff;
+
+  #pragma acc serial
+  {
+    #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 serial loop construct.  */
+
+void acc_serial_loop (void)
+{
+  int i, j, k, l, sum, diff;
+
+  #pragma acc serial 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)
+      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 now the outermost reduction clause is on
+   the serial region, not the outermost loop.  */
+
+void acc_serial_reduction (void)
+{
+  int i, j, k, sum, diff;
+
+  #pragma acc serial 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 serial loop construct, and
+   the outermost reduction clause is on that one, not the outermost loop.  */
+
+void acc_serial_loop_reduction (void)
+{
+  int i, j, k, sum, diff;
+
+  #pragma acc serial 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)
+      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 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)
+      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;
+      }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions.c
deleted file mode 100644
index bff6652..0000000
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions.c
+++ /dev/null
@@ -1,420 +0,0 @@
-/* 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;
-  }
-}
-- 
2.8.1


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