[PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message

[Frederik Harwath <frederik@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 1257 bytes --]


Hi,
This patch 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 and which hence are not orphaned.

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").  Remember that 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.

The patch has been tested by running the GCC and libgomp testsuites.
The latter tests ran with offloading to nvptx although that should not
be important here unless there was some very subtle reason for
forbidding the gang reductions on kernels loops. As expect, there seems
to be no such reason, i.e. I observed no regressions with the patch.

Can I include the patch in OG10?

Best regards,
Frederik

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

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-libgomp-Fortran-Fix-OpenACC-gang-reduction-on-an-orp.patch --]
[-- Type: text/x-diff, Size: 6361 bytes --]

From 7320635211fff3a773beb0de1914dbfcc317ab37 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
Date: Tue, 7 Jul 2020 10:41:21 +0200
Subject: [PATCH] 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-07  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 error message is not emitted for
	non-orphaned loops.

	* c-c++-common/goacc/orphan-reductions-2.c: Likewise for C and C++.
---
 gcc/fortran/openmp.c                          | 13 +++-
 .../c-c++-common/goacc/orphan-reductions-2.c  | 69 +++++++++++++++++++
 .../gfortran.dg/goacc/orphan-reductions-2.f90 | 58 ++++++++++++++++
 3 files changed, 137 insertions(+), 3 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 28408c4c99a..83c498112a8 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5926,9 +5926,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
@@ -6222,7 +6229,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/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..2b651fd2b9f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
@@ -0,0 +1,69 @@
+/* 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
+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..13887a059fe
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
@@ -0,0 +1,58 @@
+! 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 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
-- 
2.17.1