[gcc r12-5620] Re OpenACC "gang reduction on an orphan loop" error message

[gcc r12-5620] Re OpenACC "gang reduction on an orphan loop" error message

[Thomas Schwinge]

https://gcc.gnu.org/g:c4f4c60457d1657cbd72015de3d818eb6462a0e9

commit r12-5620-gc4f4c60457d1657cbd72015de3d818eb6462a0e9
Author: Frederik Harwath <frederik@codesourcery.com>
Date:   Mon Jul 20 11:24:21 2020 +0200

    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>

Diff:
---
 gcc/fortran/openmp.c                               |   9 +-
 .../c-c++-common/goacc/orphan-reductions-3.c       | 102 +++++++++++++++++++++
 .../gfortran.dg/goacc/orphan-reductions-3.f90      |  89 ++++++++++++++++++
 3 files changed, 196 insertions(+), 4 deletions(-)

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index b4100577e51..7950c7fb43d 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 00000000000..cd8ad274ebb
--- /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 00000000000..1e0b1d64578
--- /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