Re OpenACC "gang reduction on an orphan loop" error message

Follow-up to preceding commit 2b7dac2c0dcb087da9e4018943c023c0678234a3
"Make OpenACC orphan gang reductions errors".

	gcc/fortran/
	* openmp.c (oacc_is_parallel_or_serial): Evolve into...
	(oacc_is_compute_construct): ... this function.
	(resolve_oacc_loop_blocks): Use "oacc_is_compute_construct"
	instead of "oacc_is_parallel_or_serial" for checking that a
	loop is not orphaned.
	gcc/testsuite/
	* gfortran.dg/goacc/orphan-reductions-3.f90: New test
	verifying that the "gang reduction on an orphan loop" error message
	is not emitted for non-orphaned loops.
	* c-c++-common/goacc/orphan-reductions-3.c: Likewise for C and C++.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>

3 files changed, 196 insertions, 4 deletions

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index b410057..7950c7f 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -8341,9 +8341,11 @@ oacc_is_serial (gfc_code *code)
 }
 
 static bool
-oacc_is_parallel_or_serial (gfc_code *code)
+oacc_is_compute_construct (gfc_code *code)
 {
-  return oacc_is_parallel (code) || oacc_is_serial (code);
+  return (oacc_is_parallel (code)
+	  || oacc_is_kernels (code)
+	  || oacc_is_serial (code));
 }
 
 static gfc_statement
@@ -8656,8 +8658,7 @@ resolve_oacc_loop_blocks (gfc_code *code)
       for (c = omp_current_ctx; c; c = c->previous)
 	if (!oacc_is_loop (c->code))
 	  break;
-      if (c == NULL || !(oacc_is_parallel_or_serial (c->code)
-			 || oacc_is_kernels (c->code)))
+      if (c == NULL || !(oacc_is_compute_construct (c->code)))
 	gfc_error ("gang reduction on an orphan loop at %L", &code->loc);
     }
 
diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-3.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-3.c
new file mode 100644
index 0000000..cd8ad27
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-3.c
@@ -0,0 +1,102 @@
+/* Verify that the error message for gang reduction on orphaned OpenACC loops
+   is not reported for non-orphaned loops. */
+
+/* { dg-additional-options "-Wopenacc-parallelism" } */
+
+int
+kernels (int n)
+{
+  int i, s1 = 0, s2 = 0;
+#pragma acc kernels
+  {
+#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s1 = s1 + 2;
+
+#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s2 = s2 + 2;
+  }
+  return s1 + s2;
+}
+
+int
+parallel (int n)
+{
+  int i, s1 = 0, s2 = 0;
+#pragma acc parallel
+  {
+#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s1 = s1 + 2;
+
+#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s2 = s2 + 2;
+  }
+  return s1 + s2;
+}
+
+int
+serial (int n)
+{
+  int i, s1 = 0, s2 = 0;
+#pragma acc serial /* { dg-warning "region contains gang partitioned code but is not gang partitioned" } */
+  {
+#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s1 = s1 + 2;
+
+#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s2 = s2 + 2;
+  }
+  return s1 + s2;
+}
+
+int
+serial_combined (int n)
+{
+  int i, s1 = 0, s2 = 0;
+#pragma acc serial loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */
+  /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */
+  for (i = 0; i < n; i++)
+    s1 = s1 + 2;
+
+#pragma acc serial loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */
+  /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */
+  for (i = 0; i < n; i++)
+    s2 = s2 + 2;
+
+  return s1 + s2;
+}
+
+int
+parallel_combined (int n)
+{
+  int i, s1 = 0, s2 = 0;
+#pragma acc parallel loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s1 = s1 + 2;
+
+#pragma acc parallel loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s2 = s2 + 2;
+
+  return s1 + s2;
+}
+
+int
+kernels_combined (int n)
+{
+  int i, s1 = 0, s2 = 0;
+#pragma acc kernels loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s1 = s1 + 2;
+
+#pragma acc kernels loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */
+  for (i = 0; i < n; i++)
+    s2 = s2 + 2;
+
+  return s1 + s2;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-3.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-3.f90
new file mode 100644
index 0000000..1e0b1d6
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-3.f90
@@ -0,0 +1,89 @@
+! Verify that the error message for gang reductions on orphaned OpenACC loops
+! is not reported for non-orphaned loops.
+
+! { dg-additional-options "-Wopenacc-parallelism" }
+
+subroutine kernels
+  implicit none
+
+  integer, parameter :: n = 100
+  integer :: i, sum
+  sum = 0
+
+  !$acc kernels
+  !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" }
+  do i = 1, n
+     sum = sum + 1
+  end do
+  !$acc end kernels
+end subroutine kernels
+
+subroutine parallel
+  implicit none
+
+  integer, parameter :: n = 100
+  integer :: i, sum
+  sum = 0
+
+  !$acc parallel
+  !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" }
+  do i = 1, n
+     sum = sum + 1
+  end do
+  !$acc end parallel
+end subroutine parallel
+
+subroutine serial
+  implicit none
+
+  integer, parameter :: n = 100
+  integer :: i, sum
+  sum = 0
+
+  !$acc serial ! { dg-warning "region contains gang partitioned code but is not gang partitioned" }
+  !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" }
+  do i = 1, n
+     sum = sum + 1
+  end do
+  !$acc end serial
+end subroutine serial
+
+subroutine kernels_combined
+  implicit none
+
+  integer, parameter :: n = 100
+  integer :: i, sum
+  sum = 0
+
+  !$acc kernels loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" }
+  do i = 1, n
+     sum = sum + 1
+  end do
+end subroutine kernels_combined
+
+subroutine parallel_combined
+  implicit none
+
+  integer, parameter :: n = 100
+  integer :: i, sum
+  sum = 0
+
+  !$acc parallel loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" }
+  do i = 1, n
+     sum = sum + 1
+  end do
+end subroutine parallel_combined
+
+subroutine serial_combined
+  implicit none
+
+  integer, parameter :: n = 100
+  integer :: i, sum
+  sum = 0
+
+  !$acc serial loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" }
+  ! { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 }
+  do i = 1, n
+     sum = sum + 1
+  end do
+end subroutine serial_combined