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

[Thomas Schwinge]

Hi!

On 2020-11-03T09:39:54+0100, I wrote:
> On 2018-12-20T15:28:33+0100, I wrote:
>> On behalf of Gergő (who doesn't have write access yet) I've pushed the
>> attached to openacc-gcc-8-branch.
>
> (Which then eventually got into master branch via Frederik.)

>> --- /dev/null
>> +++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
>> @@ -0,0 +1,492 @@
>> +/* Test erroneous cases of nested reduction loops.  */
>> +
>> +void acc_parallel (void)
>> +{
>> +  int i, j, k, l, sum, diff;
>> +
>> +  #pragma acc parallel
>> +  {
>> +    #pragma acc loop reduction(+:sum)
>> +    for (i = 0; i < 10; i++)
>> +      #pragma acc loop // { dg-error "nested loop in reduction needs 
>> reduction clause for .sum." }
>> +      for (j = 0; j < 10; j++)
>> +        #pragma acc loop reduction(+:sum)
>> +        for (k = 0; k < 10; k++)
>> +          sum = 1;
>
>> +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;

Getting these diagnostics consistent is easy enough, however.  I've
pushed "[OpenACC] Enable inconsistent nested 'reduction' clauses checking
for OpenACC 'kernels'" to master branch in commit
64dc14b1a764bd3059170431c9b43c6192dbd48f, and backported to
releases/gcc-10 branch in commit
217fb4d4e59e7d6e03a3704f80f401e2a641dbe5, see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander 
Walter

>From 64dc14b1a764bd3059170431c9b43c6192dbd48f Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Thu, 22 Oct 2020 11:04:22 +0200
Subject: [PATCH] [OpenACC] Enable inconsistent nested 'reduction' clauses
 checking for OpenACC 'kernels'

	gcc/
	* omp-low.c (scan_omp_for) <OpenACC>: Move earlier inconsistent
	nested 'reduction' clauses checking.
	gcc/testsuite/
	* c-c++-common/goacc/nested-reductions-1-kernels.c: Extend.
	* c-c++-common/goacc/nested-reductions-2-kernels.c: Likewise.
	* gfortran.dg/goacc/nested-reductions-1-kernels.f90: Likewise.
	* gfortran.dg/goacc/nested-reductions-2-kernels.f90: Likewise.
---
 gcc/omp-low.c                                 |  36 +-
 .../goacc/nested-reductions-1-kernels.c       | 199 +++++++++-
 .../goacc/nested-reductions-2-kernels.c       | 271 +++++++++++++-
 .../goacc/nested-reductions-1-kernels.f90     | 251 ++++++++++++-
 .../goacc/nested-reductions-2-kernels.f90     | 346 +++++++++++++++++-
 5 files changed, 1063 insertions(+), 40 deletions(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index de5142f979b0..2f1a544bd46e 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2454,23 +2454,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
 	  }
 
       if (tgt && is_oacc_kernels (tgt))
-	{
-	  /* Strip out reductions, as they are not handled yet.  */
-	  tree *prev_ptr = &clauses;
-
-	  while (tree probe = *prev_ptr)
-	    {
-	      tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
-	      
-	      if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
-		*prev_ptr = *next_ptr;
-	      else
-		prev_ptr = next_ptr;
-	    }
-
-	  gimple_omp_for_set_clauses (stmt, clauses);
-	  check_oacc_kernel_gwv (stmt, ctx);
-	}
+	check_oacc_kernel_gwv (stmt, ctx);
 
       /* Collect all variables named in reductions on this loop.  Ensure
 	 that, if this loop has a reduction on some variable v, and there is
@@ -2553,6 +2537,24 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
       ctx->outer_reduction_clauses
 	= chainon (unshare_expr (ctx->local_reduction_clauses),
 		   ctx->outer_reduction_clauses);
+
+      if (tgt && is_oacc_kernels (tgt))
+	{
+	  /* Strip out reductions, as they are not handled yet.  */
+	  tree *prev_ptr = &clauses;
+
+	  while (tree probe = *prev_ptr)
+	    {
+	      tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
+
+	      if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
+		*prev_ptr = *next_ptr;
+	      else
+		prev_ptr = next_ptr;
+	    }
+
+	  gimple_omp_for_set_clauses (stmt, clauses);
+	}
     }
 
   scan_sharing_clauses (clauses, ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
index 68cb8f82ee57..9323e2c8d7e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
@@ -6,8 +6,6 @@ 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)
@@ -16,6 +14,12 @@ void acc_kernels (void)
         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)
@@ -25,17 +29,208 @@ void acc_kernels (void)
 
     #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) // { dg-warning "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-warning "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-warning "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;
+      }
+  }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+   the kernels region, not the outermost loop.  */
+
+void acc_kernels_reduction (void)
+{
+  /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+     supported on the 'kernels' construct.  */
+}
+
+/* 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) // { dg-warning "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-warning "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-warning "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;
+      }
+
+    #pragma acc loop reduction(+:sum)
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(+:sum) // { dg-warning "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-warning "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;
+      }
+
+    #pragma acc loop reduction(+:sum)
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(+:sum) // { dg-warning "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 // { dg-warning "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-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
index 80d7c53daca7..dec7dbda45dc 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
@@ -4,36 +4,107 @@
 
 void acc_kernels (void)
 {
-  int i, j, k, sum, diff;
+  int i, j, k, l, 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++)
+      #pragma acc loop // { dg-warning "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
+      #pragma acc loop collapse(2) // { dg-warning "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-warning "nested loop in reduction needs reduction clause for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "conflicting reduction operations for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "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(-:diff)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
       for (j = 0; j < 10; j++)
-        #pragma acc loop
+        #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
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "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-warning "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 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 // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
       for (j = 0; j < 10; j++)
         #pragma acc loop reduction(+:sum)
         for (k = 0; k < 10; k++)
@@ -41,10 +112,194 @@ void acc_kernels (void)
 
     #pragma acc loop reduction(+:sum)
     for (i = 0; i < 10; i++)
-      #pragma acc loop reduction(-:sum)
+      #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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 "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "nested loop in reduction needs reduction clause for .diff." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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 kernels region, not the outermost loop.  */
+void acc_kernels_reduction (void)
+{
+  /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+     supported on the 'kernels' construct.  */
+}
+
+/* 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 // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
       for (j = 0; j < 10; j++)
         #pragma acc loop reduction(+:sum)
         for (k = 0; k < 10; k++)
           sum = 1;
+
+    #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+          #pragma acc loop reduction(+:sum)
+          for (l = 0; l < 10; l++)
+            sum = 1;
+
+    #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+        for (k = 0; k < 10; k++)
+          sum = 1;
+
+    #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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) // { dg-warning "conflicting reduction operations for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+	  for (l = 0; l < 10; l++)
+	    sum = 1;
+
+    #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+      #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." })
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+	  for (l = 0; l < 10; l++)
+	    sum = 1;
+
+    #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "nested loop in reduction needs reduction clause for .diff." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (j = 0; j < 10; j++)
+          #pragma acc loop reduction(-:diff)
+          for (k = 0; k < 10; k++)
+            diff = 1;
+      }
   }
 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
index 0999f8767c6e..60cb63006c3c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
@@ -6,8 +6,6 @@ subroutine acc_kernels ()
   implicit none (type, external)
   integer :: i, j, k, sum, diff
 
-  ! FIXME:  These tests are not meaningful yet because reductions in
-  ! kernels regions are not supported yet.
   !$acc kernels
     !$acc loop reduction(+:sum)
     do i = 1, 10
@@ -18,6 +16,15 @@ subroutine acc_kernels ()
       end do
     end do
 
+    !$acc loop collapse(2) reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
     !$acc loop reduction(+:sum)
     do i = 1, 10
       !$acc loop reduction(+:sum)
@@ -28,6 +35,15 @@ subroutine acc_kernels ()
       end do
     end do
 
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop collapse(2) reduction(+:sum)
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
@@ -39,7 +55,6 @@ subroutine acc_kernels ()
       end do
     end do
 
-
     !$acc loop reduction(+:sum)
     do i = 1, 10
       !$acc loop reduction(+:sum)
@@ -51,5 +66,235 @@ subroutine acc_kernels ()
       end do
     end do
 
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
   !$acc end kernels
 end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+  implicit none (type, external)
+  integer :: h, i, j, k, l, sum, diff
+
+  !$acc kernels loop
+  do h = 1, 10
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop collapse(2) reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop collapse(2) reduction(+:sum)
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop.  */
+
+subroutine acc_kernels_reduction ()
+  implicit none (type, external)
+
+  ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+  ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! 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.  */
+subroutine acc_kernels_loop_reduction ()
+  implicit none (type, external)
+  integer :: h, i, j, k, sum, diff
+
+  !$acc kernels loop reduction(+:sum)
+  do h = 1, 10
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    do i = 1, 10
+      !$acc loop
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    do i = 1, 10
+      do j = 1, 10
+        !$acc loop
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop_reduction
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
index edfd8625faf3..6ee41843d31a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
@@ -3,14 +3,15 @@
 ! See also 'c-c++-common/goacc/nested-reductions-2-kernels.c'.
 
 subroutine acc_kernels ()
-  integer :: i, j, k, sum, diff
+  implicit none (type, external)
+  integer :: i, j, k, l, sum, diff
 
-  ! FIXME:  No diagnostics are produced for these loops because reductions
-  ! in kernels regions are not supported yet.
   !$acc kernels
     !$acc loop reduction(+:sum)
     do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
       do j = 1, 10
+        !$acc loop reduction(+:sum)
         do k = 1, 10
           sum = 1
         end do
@@ -19,19 +20,37 @@ subroutine acc_kernels ()
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop
+      !$acc loop collapse(2)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
       do j = 1, 10
         do k = 1, 10
-          sum = 1
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
         end do
       end do
     end do
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop reduction(-:diff)
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
       do j = 1, 10
-        !$acc loop
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
         do k = 1, 10
           sum = 1
         end do
@@ -40,9 +59,9 @@ subroutine acc_kernels ()
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
       do j = 1, 10
-        !$acc loop reduction(+:sum)
+        !$acc loop reduction(-:sum)
         do k = 1, 10
           sum = 1
         end do
@@ -51,13 +70,320 @@ subroutine acc_kernels ()
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop reduction(-:sum)
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
       do j = 1, 10
         !$acc loop reduction(+:sum)
         do k = 1, 10
           sum = 1
         end do
       end do
+
+      !$acc loop reduction(+:sum)  ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
     end do
   !$acc end kernels
 end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+  implicit none (type, external)
+  integer :: h, i, j, k, l, sum, diff
+
+  !$acc kernels loop
+  do h = 1, 10
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop collapse(2)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(+:sum)  ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop.
+
+subroutine acc_kernels_reduction ()
+  implicit none (type, external)
+
+  ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+  ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! 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.  */
+subroutine acc_kernels_loop_reduction ()
+  implicit none (type, external)
+  integer :: h, i, j, k, l, sum, diff
+
+  !$acc kernels loop reduction(+:sum)
+  do h = 1, 10
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop collapse(2)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }  
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(max:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(max:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(+:sum)  ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop_reduction
-- 
2.17.1


>From 217fb4d4e59e7d6e03a3704f80f401e2a641dbe5 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Thu, 22 Oct 2020 11:04:22 +0200
Subject: [PATCH] [OpenACC] Enable inconsistent nested 'reduction' clauses
 checking for OpenACC 'kernels'

	gcc/
	* omp-low.c (scan_omp_for) <OpenACC>: Move earlier inconsistent
	nested 'reduction' clauses checking.
	gcc/testsuite/
	* c-c++-common/goacc/nested-reductions-1-kernels.c: Extend.
	* c-c++-common/goacc/nested-reductions-2-kernels.c: Likewise.
	* gfortran.dg/goacc/nested-reductions-1-kernels.f90: Likewise.
	* gfortran.dg/goacc/nested-reductions-2-kernels.f90: Likewise.

(cherry picked from commit 64dc14b1a764bd3059170431c9b43c6192dbd48f)
---
 gcc/omp-low.c                                 |  36 +-
 .../goacc/nested-reductions-1-kernels.c       | 199 +++++++++-
 .../goacc/nested-reductions-2-kernels.c       | 271 +++++++++++++-
 .../goacc/nested-reductions-1-kernels.f90     | 251 ++++++++++++-
 .../goacc/nested-reductions-2-kernels.f90     | 346 +++++++++++++++++-
 5 files changed, 1063 insertions(+), 40 deletions(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 16afc3720af2..91e1bcc6b69a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2447,23 +2447,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
 	  }
 
       if (tgt && is_oacc_kernels (tgt))
-	{
-	  /* Strip out reductions, as they are not handled yet.  */
-	  tree *prev_ptr = &clauses;
-
-	  while (tree probe = *prev_ptr)
-	    {
-	      tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
-	      
-	      if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
-		*prev_ptr = *next_ptr;
-	      else
-		prev_ptr = next_ptr;
-	    }
-
-	  gimple_omp_for_set_clauses (stmt, clauses);
-	  check_oacc_kernel_gwv (stmt, ctx);
-	}
+	check_oacc_kernel_gwv (stmt, ctx);
 
       /* Collect all variables named in reductions on this loop.  Ensure
 	 that, if this loop has a reduction on some variable v, and there is
@@ -2546,6 +2530,24 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
       ctx->outer_reduction_clauses
 	= chainon (unshare_expr (ctx->local_reduction_clauses),
 		   ctx->outer_reduction_clauses);
+
+      if (tgt && is_oacc_kernels (tgt))
+	{
+	  /* Strip out reductions, as they are not handled yet.  */
+	  tree *prev_ptr = &clauses;
+
+	  while (tree probe = *prev_ptr)
+	    {
+	      tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
+
+	      if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
+		*prev_ptr = *next_ptr;
+	      else
+		prev_ptr = next_ptr;
+	    }
+
+	  gimple_omp_for_set_clauses (stmt, clauses);
+	}
     }
 
   scan_sharing_clauses (clauses, ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
index 68cb8f82ee57..9323e2c8d7e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
@@ -6,8 +6,6 @@ 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)
@@ -16,6 +14,12 @@ void acc_kernels (void)
         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)
@@ -25,17 +29,208 @@ void acc_kernels (void)
 
     #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) // { dg-warning "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-warning "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-warning "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;
+      }
+  }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+   the kernels region, not the outermost loop.  */
+
+void acc_kernels_reduction (void)
+{
+  /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+     supported on the 'kernels' construct.  */
+}
+
+/* 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) // { dg-warning "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-warning "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-warning "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;
+      }
+
+    #pragma acc loop reduction(+:sum)
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(+:sum) // { dg-warning "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-warning "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;
+      }
+
+    #pragma acc loop reduction(+:sum)
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(+:sum) // { dg-warning "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 // { dg-warning "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-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
index 80d7c53daca7..dec7dbda45dc 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
@@ -4,36 +4,107 @@
 
 void acc_kernels (void)
 {
-  int i, j, k, sum, diff;
+  int i, j, k, l, 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++)
+      #pragma acc loop // { dg-warning "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
+      #pragma acc loop collapse(2) // { dg-warning "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-warning "nested loop in reduction needs reduction clause for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "conflicting reduction operations for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "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(-:diff)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
       for (j = 0; j < 10; j++)
-        #pragma acc loop
+        #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
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." }
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "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-warning "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 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 // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
       for (j = 0; j < 10; j++)
         #pragma acc loop reduction(+:sum)
         for (k = 0; k < 10; k++)
@@ -41,10 +112,194 @@ void acc_kernels (void)
 
     #pragma acc loop reduction(+:sum)
     for (i = 0; i < 10; i++)
-      #pragma acc loop reduction(-:sum)
+      #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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 "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "nested loop in reduction needs reduction clause for .diff." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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 kernels region, not the outermost loop.  */
+void acc_kernels_reduction (void)
+{
+  /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+     supported on the 'kernels' construct.  */
+}
+
+/* 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 // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
       for (j = 0; j < 10; j++)
         #pragma acc loop reduction(+:sum)
         for (k = 0; k < 10; k++)
           sum = 1;
+
+    #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+          #pragma acc loop reduction(+:sum)
+          for (l = 0; l < 10; l++)
+            sum = 1;
+
+    #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+        for (k = 0; k < 10; k++)
+          sum = 1;
+
+    #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      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) // { dg-warning "conflicting reduction operations for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+        #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+	  for (l = 0; l < 10; l++)
+	    sum = 1;
+
+    #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." }
+    for (i = 0; i < 10; i++)
+      #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+      for (j = 0; j < 10; j++)
+      #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." })
+      // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (k = 0; k < 10; k++)
+	  #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+	  for (l = 0; l < 10; l++)
+	    sum = 1;
+
+    #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    for (i = 0; i < 10; i++)
+      {
+        #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        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-warning "nested loop in reduction needs reduction clause for .diff." }
+        // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } 
+        for (j = 0; j < 10; j++)
+          #pragma acc loop reduction(-:diff)
+          for (k = 0; k < 10; k++)
+            diff = 1;
+      }
   }
 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
index 0999f8767c6e..60cb63006c3c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
@@ -6,8 +6,6 @@ subroutine acc_kernels ()
   implicit none (type, external)
   integer :: i, j, k, sum, diff
 
-  ! FIXME:  These tests are not meaningful yet because reductions in
-  ! kernels regions are not supported yet.
   !$acc kernels
     !$acc loop reduction(+:sum)
     do i = 1, 10
@@ -18,6 +16,15 @@ subroutine acc_kernels ()
       end do
     end do
 
+    !$acc loop collapse(2) reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
     !$acc loop reduction(+:sum)
     do i = 1, 10
       !$acc loop reduction(+:sum)
@@ -28,6 +35,15 @@ subroutine acc_kernels ()
       end do
     end do
 
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop collapse(2) reduction(+:sum)
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
@@ -39,7 +55,6 @@ subroutine acc_kernels ()
       end do
     end do
 
-
     !$acc loop reduction(+:sum)
     do i = 1, 10
       !$acc loop reduction(+:sum)
@@ -51,5 +66,235 @@ subroutine acc_kernels ()
       end do
     end do
 
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
   !$acc end kernels
 end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+  implicit none (type, external)
+  integer :: h, i, j, k, l, sum, diff
+
+  !$acc kernels loop
+  do h = 1, 10
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop collapse(2) reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop collapse(2) reduction(+:sum)
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop.  */
+
+subroutine acc_kernels_reduction ()
+  implicit none (type, external)
+
+  ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+  ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! 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.  */
+subroutine acc_kernels_loop_reduction ()
+  implicit none (type, external)
+  integer :: h, i, j, k, sum, diff
+
+  !$acc kernels loop reduction(+:sum)
+  do h = 1, 10
+    do i = 1, 10
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    do i = 1, 10
+      !$acc loop
+      do j = 1, 10
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    do i = 1, 10
+      do j = 1, 10
+        !$acc loop
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(-:diff)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(+:sum)  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop  ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop_reduction
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
index edfd8625faf3..6ee41843d31a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
@@ -3,14 +3,15 @@
 ! See also 'c-c++-common/goacc/nested-reductions-2-kernels.c'.
 
 subroutine acc_kernels ()
-  integer :: i, j, k, sum, diff
+  implicit none (type, external)
+  integer :: i, j, k, l, sum, diff
 
-  ! FIXME:  No diagnostics are produced for these loops because reductions
-  ! in kernels regions are not supported yet.
   !$acc kernels
     !$acc loop reduction(+:sum)
     do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
       do j = 1, 10
+        !$acc loop reduction(+:sum)
         do k = 1, 10
           sum = 1
         end do
@@ -19,19 +20,37 @@ subroutine acc_kernels ()
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop
+      !$acc loop collapse(2)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
       do j = 1, 10
         do k = 1, 10
-          sum = 1
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
         end do
       end do
     end do
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop reduction(-:diff)
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
       do j = 1, 10
-        !$acc loop
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
         do k = 1, 10
           sum = 1
         end do
@@ -40,9 +59,9 @@ subroutine acc_kernels ()
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
       do j = 1, 10
-        !$acc loop reduction(+:sum)
+        !$acc loop reduction(-:sum)
         do k = 1, 10
           sum = 1
         end do
@@ -51,13 +70,320 @@ subroutine acc_kernels ()
 
     !$acc loop reduction(+:sum)
     do i = 1, 10
-      !$acc loop reduction(-:sum)
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
       do j = 1, 10
         !$acc loop reduction(+:sum)
         do k = 1, 10
           sum = 1
         end do
       end do
+
+      !$acc loop reduction(+:sum)  ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
     end do
   !$acc end kernels
 end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+  implicit none (type, external)
+  integer :: h, i, j, k, l, sum, diff
+
+  !$acc kernels loop
+  do h = 1, 10
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop collapse(2)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum)
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(+:sum) reduction(-:diff)
+    do i = 1, 10
+      !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(+:sum)  ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop.
+
+subroutine acc_kernels_reduction ()
+  implicit none (type, external)
+
+  ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+  ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! 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.  */
+subroutine acc_kernels_loop_reduction ()
+  implicit none (type, external)
+  integer :: h, i, j, k, l, sum, diff
+
+  !$acc kernels loop reduction(+:sum)
+  do h = 1, 10
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop collapse(2)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(+:sum)
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }  
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(max:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(max:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+        do k = 1, 10
+          !$acc loop reduction(*:sum)  ! { dg-warning "conflicting reduction operations for .sum." }
+          do l = 1, 10
+            sum = 1
+          end do
+        end do
+      end do
+    end do
+
+    !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+    do i = 1, 10
+      !$acc loop reduction(-:diff)  ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(+:sum)
+        do k = 1, 10
+          sum = 1
+        end do
+      end do
+
+      !$acc loop reduction(+:sum)  ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+      ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+      do j = 1, 10
+        !$acc loop reduction(-:diff)
+        do k = 1, 10
+          diff = 1
+        end do
+      end do
+    end do
+  end do
+end subroutine acc_kernels_loop_reduction
-- 
2.17.1