diff mbox series

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

Message ID 07276106-5238-482c-31df-41a90f389783@mentor.com
State New
Headers show
Series [og8] Update code and reduction tests for `serial' construct | expand

Commit Message

Gergö Barany Dec. 21, 2018, 12:24 p.m. UTC
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.

Comments

Thomas Schwinge Dec. 21, 2018, 1:29 p.m. UTC | #1
Hi Gergő!

On Fri, 21 Dec 2018 13:24:34 +0100, Gergö Barany <gergo@codesourcery.com> wrote:
> 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.

My bad for not noticing that one before pushing these changes
yesterday...  ;-|

> 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?

Yes, thanks.  To record the review effort, please include "Reviewed-by:
Thomas Schwinge <thomas@codesourcery.com>" in the commit log, see
<https://gcc.gnu.org/wiki/Reviewed-by>.

Just two minor notes:

> 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

Might want to say "OpenACC `serial' construct" here, so that it's
obviously clear what this is about.

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

The paths used inside the snippets are always relative to the respective
ChangeLog file, so here:

> 	gcc/testsuite/
> 	* c-c++-common/goacc/nested-reductions-fail.c: Renamed to...
> 	[...]

You got that right in the "gcc/testsuite/ChangeLog.openacc" file, but it
needs to be updated in the commit message.  (... which in GCC typically
is just a copy'n'paste of the respective ChangeLog updates, possibly with
some introductory text to describe the "why" of the commit etc., instead
of just the "how" as done by the GNU ChangeLogs.)  For avoidance of
doubt, in your example here, I'd say that the summary line plus ChangeLog
is sufficient.


Grüße
 Thomas


> ---
>  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
>
diff mbox series

Patch

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