[gcc/devel/omp/gcc-11] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message

[gcc/devel/omp/gcc-11] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message

[Kwok Yeung]

https://gcc.gnu.org/g:99e3e66aa087efd1154ad6762e74d348d20f2784

commit 99e3e66aa087efd1154ad6762e74d348d20f2784
Author: Frederik Harwath <frederik@codesourcery.com>
Date:   Mon Jul 20 11:24:21 2020 +0200

    libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message
    
    According to the OpenACC standard version 2.5 and later, reductions on
    orphaned gang loops are explicitly disallowed (cf.  section "Changes
    from Version 2.0 to 2.5").  A loop is "orphaned" if it is not
    lexically contained in a compute construct (cf. section "Loop
    construct" of the OpenACC standard), i.e. in either a "parallel", a
    "serial", or a "kernels" construct.
    
    This commit fixes the check for reductions on orphaned gang loops in
    the Fortran frontend which (in contrast to the C, C++ frontends)
    erroneously rejects reductions on gang loops that are contained in
    "kernels" constructs.
    
    2020-07-20  Frederik Harwath  <frederik@codesourcery.com>
    
    gcc/fortran/
    
            * openmp.c (oacc_is_parallel_or_serial): Removed function.
            (oacc_is_kernels): New function.
            (oacc_is_compute_construct): New 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-2.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-2.c: Likewise for C and C++.

Diff:
---
 gcc/fortran/ChangeLog.omp                          |   9 ++
 gcc/fortran/openmp.c                               |  13 ++-
 gcc/testsuite/ChangeLog.omp                        |   7 ++
 .../c-c++-common/goacc/orphan-reductions-2.c       | 103 +++++++++++++++++++++
 .../gfortran.dg/goacc/orphan-reductions-2.f90      |  87 +++++++++++++++++
 5 files changed, 216 insertions(+), 3 deletions(-)

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 66897ec2f20..e619434815e 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,12 @@
+2020-07-20  Frederik Harwath  <frederik@codesourcery.com>
+
+	* openmp.c (oacc_is_parallel_or_serial): Removed function.
+	(oacc_is_kernels): New function.
+	(oacc_is_compute_construct): New 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.
+
 2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
 	    Gergö Barany <gergo@codesourcery.com>
 
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 5421af669f8..7b344aa8223 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -6700,9 +6700,16 @@ oacc_is_serial (gfc_code *code)
 }
 
 static bool
-oacc_is_parallel_or_serial (gfc_code *code)
+oacc_is_kernels (gfc_code *code)
 {
-  return oacc_is_parallel (code) || oacc_is_serial (code);
+  return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP;
+}
+
+static bool
+oacc_is_compute_construct (gfc_code *code)
+{
+  return oacc_is_parallel (code) || oacc_is_serial (code)
+    || oacc_is_kernels (code);
 }
 
 static gfc_statement
@@ -6978,7 +6985,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))
+      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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index a8a41238c4e..87044b86e3c 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,10 @@
+2020-07-20  Frederik Harwath  <frederik@codesourcery.com>
+
+	* gfortran.dg/goacc/orphan-reductions-2.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-2.c: Likewise for C and C++.
+
 2020-06-02  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
 	* c-c++-common/goacc/noncontig_array-1.c: Dump Gimple pass.
diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
new file mode 100644
index 00000000000..d30321710dd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
@@ -0,0 +1,103 @@
+/* Verify that the error message for gang reduction on orphaned OpenACC loops
+   is not reported for non-orphaned loops. */
+
+#include <assert.h>
+
+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-2.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
new file mode 100644
index 00000000000..6ad38039696
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
@@ -0,0 +1,87 @@
+! Verify that the error message for gang reductions on orphaned OpenACC loops
+! is not reported for non-orphaned loops.
+
+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