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

Message ID xu8f8sfv4shz.fsf@harwath.name
State New
Headers show
Series
  • [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message
Related show

Commit Message

Frederik Harwath July 7, 2020, 8:52 a.m.
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

Comments

Thomas Schwinge July 7, 2020, 10:42 a.m. | #1
Hi Frederik!

(CC <fortran@gcc.gnu.org> added, for everything touching gfortran.)

On 2020-07-07T10:52:08+0200, Frederik Harwath <frederik@codesourcery.com> wrote:
> This patch fixes the check for reductions on orphaned gang loops


This is the "Make OpenACC orphan gang reductions errors" functionality
originally added in gomp-4_0-branch r247461.

> 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.


Or the other way round: a 'loop' construct is orphaned if it appears
inside a 'routine' region, right?

> 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.


Note that the aforementioned gomp-4_0-branch r247461,
openacc-gcc-7-branch commit 0554f9f79325960c72166327d442a553cd35bad9, and
openacc-gcc-8-branch commit 65dd9cf3b3c45d64d72967df1e4a54778cb4e35f
still do contain the appropriate 'kernels' handling.  Just in
openacc-gcc-9-branch commit 533beb2ec19f8486e4b1b645a153746f96b41f04 this
got (a) mixed together with a bunch of other, unrelated changes ("Various
OpenACC reduction enhancements"), and (b) the 'kernels' handling got
removed.  Julian (Git author), or Kwok (Git committer), do you remember
any rationale for that?  Later, this then got picked into devel/omp/gcc-9
commit 3fa4bb72dcb3b9171952a0eca5310bb8811d5ffd, and devel/omp/gcc-10
commit 6b3e1f7f05cd360bbd356b3f78511aa2ec3f40c3.

> Can I include the patch in OG10?


Unless Julian/Kwok speak up soon: OK, thanks.

    Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>


May want to remove "libgomp" from the first line of the commit log --
this commit doesn't relate to libgomp specifically.

(Ideally, we'd also test 'serial' construct in addition to 'kernels',
'parallel', but we can add that later.  I anyway have a WIP patch
waiting, adding more 'serial' construct testing, for a different reason,
so I'll include it there.)


Grüße
 Thomas


> 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

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Frederik Harwath July 7, 2020, 12:39 p.m. | #2
Thomas Schwinge <thomas@codesourcery.com> writes:

Hi Thomas,

> (CC <fortran@gcc.gnu.org> added, for everything touching gfortran.)


Thanks!

> On 2020-07-07T10:52:08+0200, Frederik Harwath <frederik@codesourcery.com> wrote:

>> This patch fixes the check for reductions on orphaned gang loops

>

> This is the "Make OpenACC orphan gang reductions errors" functionality

> originally added in gomp-4_0-branch r247461.

>

>> 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.

>

> Or the other way round: a 'loop' construct is orphaned if it appears

> inside a 'routine' region, right?


The "not lexically contained in a compute construct" definition is
from the standard. Assuming that the frontend's parser rejects "loop"
directives if they do not occur inside of either the "serial",
"parallel", "kernels" compute constructs or in a function with a
"routine" directive, both definitions should be indeed equivalent ;-).

> Unless Julian/Kwok speak up soon: OK, thanks.

>

>     Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

>

> May want to remove "libgomp" from the first line of the commit log --

> this commit doesn't relate to libgomp specifically.


Right.

> (Ideally, we'd also test 'serial' construct in addition to 'kernels',

> 'parallel', but we can add that later.  I anyway have a WIP patch

> waiting, adding more 'serial' construct testing, for a different reason,

> so I'll include it there.)


I had left this out intentionally, because having the gang reduction in
the serial construct leads to a "region contains gang partitioned
code but is not gang partitioned"
error. Of course, we might still add a test case with that expectation.

Thanks for the review!

Frederik
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Frederik Harwath July 20, 2020, 10:26 a.m. | #3
Thomas Schwinge <thomas@codesourcery.com> writes:

Hi Thomas,

>> Can I include the patch in OG10?

>

> Unless Julian/Kwok speak up soon: OK, thanks.


This has been delayed a bit by my vacation, but I have now committed
the patch.

> May want to remove "libgomp" from the first line of the commit log --

> this commit doesn't relate to libgomp specifically.

>

> (Ideally, we'd also test 'serial' construct in addition to 'kernels',

> 'parallel', but we can add that later.  I anyway have a WIP patch

> waiting, adding more 'serial' construct testing, for a different reason,

> so I'll include it there.)


I forgot to remove "libgomp" from the commit message, sorry, but
I have included the test cases for the "serial construct".

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
From 7c10ae450b95495dda362cb66770bb78b546592e Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>

Date: Mon, 20 Jul 2020 11:24: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-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++.
---
 gcc/fortran/ChangeLog                         |   9 ++
 gcc/fortran/openmp.c                          |  13 ++-
 gcc/testsuite/ChangeLog                       |   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(-)
 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/ChangeLog b/gcc/fortran/ChangeLog
index e86279cb647..5a1f81c286e 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -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-07-08  Harald Anlauf  <anlauf@gmx.de>
 
 	Backported from master:
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index ab68e9f2173..706933c869a 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5927,9 +5927,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
@@ -6223,7 +6230,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 b/gcc/testsuite/ChangeLog
index 59e6c93b07a..fa1937a4ea2 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -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-07-12  Jakub Jelinek  <jakub@redhat.com>
 
 	Backported from master:
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
-- 
2.17.1

Patch

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