Hi!
On 2020-03-19T17:12:02+0000, Kwok Cheung Yeung <kwok_ye...@mentor.com> wrote:
> On 18/03/2020 11:34 pm, Kwok Cheung Yeung wrote:
>> I was looking at the regression in c-c++-common/goacc/nested-reductions.c,
>> which
>> has the following excess warnings in acc_routine:
>>
>> /scratch/kyeung/openacc/og10/nvidia/src/gcc-og10-branch/gcc/testsuite/c-c++-common/goacc/nested-reductions.c:360:15:
>> warning: insufficient partitioning available to parallelize loop
>> /scratch/kyeung/openacc/og10/nvidia/src/gcc-og10-branch/gcc/testsuite/c-c++-common/goacc/nested-reductions.c:369:17:
>> warning: insufficient partitioning available to parallelize loop
>> /scratch/kyeung/openacc/og10/nvidia/src/gcc-og10-branch/gcc/testsuite/c-c++-common/goacc/nested-reductions.c:375:17:
>> warning: insufficient partitioning available to parallelize loop
>> /scratch/kyeung/openacc/og10/nvidia/src/gcc-og10-branch/gcc/testsuite/c-c++-common/goacc/nested-reductions.c:320:6:
>> warning: region is gang partitioned but does not contain gang partitioned
>> code
>>
>> It is caused by the following code in the patch 'Make OpenACC orphan
>> gang reductions errors"] (originally by Cesar):
>>
>> + /* Orphan reductions cannot have gang partitioning. */
>> + if ((loop->flags & OLF_REDUCTION)
>> + && oacc_get_fn_attrib (current_function_decl)
>> + && !lookup_attribute ("omp target entrypoint",
>> + DECL_ATTRIBUTES (current_function_decl)))
>> + this_mask = GOMP_DIM_MASK (GOMP_DIM_WORKER);
Right. However, that code doesn't implement what the OpenACC
specification actually says. ;-)
>> The problem is that acc_routine is not declared with 'omp target entrypoint',
>> but it does have '#pragma acc_routine gang' applied to it. From what I
>> understand of the OpenACC spec, this means that the function can be called
>> from
>> the accelerator, and may contain a loop at the gang-level.
Right.
>> So is allowing gang
>> reductions for functions with '#pragma acc_routine gang' (but not for worker
>> or
>> vector) the right thing to do here?
No, that's precisely the thing that the compiler needs to diagnose. See
OpenACC 2.6, 2.9.11. "reduction clause", which places a restriction such
that "The 'reduction' clause may not be specified on an orphaned 'loop'
construct with the 'gang' clause, or on an orphaned 'loop' construct that
will generate gang parallelism in a procedure that is compiled with the
'routine gang' clause." */
Cesar apparently read the last part to mean that inside a 'routine gang',
a 'loop reduction' with implicit 'gang' level of parallelism should be
demoted to 'worker' level of parallelism. But what actually is meant,
simply, is that in such cases we raise the same "gang reduction on an
orphan loop" error diagnostic that we raise for explicit 'gang' level of
parallelism. (..., and adjust our offending test cases).
Now, re your og10 etc. change:
> Allow gang-level reductions in OpenACC routines with gang-level
> parallelism
> gcc/
> * omp-offload.c (oacc_loop_auto_partitions): Check for 'omp declare
> target' attributes with a gang clause attached.
> --- a/gcc/omp-offload.c
> +++ b/gcc/omp-offload.c
> @@ -1374,14 +1374,32 @@ oacc_loop_auto_partitions (oacc_loop *loop, unsigned
> outer_mask,
> /* Orphan reductions cannot have gang partitioning. */
> if ((loop->flags & OLF_REDUCTION)
> - && oacc_get_fn_attrib (current_function_decl)
> - && !lookup_attribute ("omp target entrypoint",
> + && oacc_get_fn_attrib (current_function_decl))
> + {
> + bool gang_p = false;
> + tree attr
> + = lookup_attribute ("omp declare target",
> + DECL_ATTRIBUTES (current_function_decl));
> +
> + if (attr)
> + for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG)
> + {
> + gang_p = true;
> + break;
> + }
> +
> + if (lookup_attribute ("omp target entrypoint",
> DECL_ATTRIBUTES (current_function_decl)))
> - this_mask = GOMP_DIM_MASK (GOMP_DIM_WORKER);
> + gang_p = true;
> +
> + if (!gang_p)
> + this_mask = GOMP_DIM_MASK (GOMP_DIM_WORKER);
> + }
..., I don't understand what exactly that is meant to do: as far as I can
tell, we always get 'gang_p == true' from that code?
Instead, I've pushed to master branch
commit 365cd5f9ba812c389b404a53d99ab5dded5097f4 '[OpenACC] Remove
erroneous "Orphan reductions cannot have gang partitioning" handling',
see attached. This implements the desired "gang reduction on an orphan
loop" error diagnostics also for these implicit 'gang' cases, via the
middle-end checking that I've just added in
commit 77d24d43644909852998043335b5a0e09d1e8f02
'Consolidate OpenACC "gang reduction on an orphan loop" checking'.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
>From 365cd5f9ba812c389b404a53d99ab5dded5097f4 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri, 26 Nov 2021 13:11:16 +0100
Subject: [PATCH] [OpenACC] Remove erroneous "Orphan reductions cannot have
gang partitioning" handling
That is:
-/* Ensure that the middle end does not assign gang level parallelism
- to orphan loop containing reductions. */
+/* Verify that we diagnose "gang reduction on an orphan loop" for automatically
+ assigned gang level of parallelism. */
... to implement what the OpenACC specification actually says.
Fix-up for preceding commit 2b7dac2c0dcb087da9e4018943c023c0678234a3
"Make OpenACC orphan gang reductions errors".
gcc/
* omp-offload.c (oacc_loop_auto_partitions): Remove erroneous
"Orphan reductions cannot have gang partitioning" handling.
gcc/testsuite/
* c-c++-common/goacc/nested-reductions-1-routine.c: Adjust.
* c-c++-common/goacc/nested-reductions-2-routine.c: Adjust.
* c-c++-common/goacc/orphan-reductions-2.c: Adjust.
* gfortran.dg/goacc/nested-reductions-1-routine.f90: Adjust.
* gfortran.dg/goacc/nested-reductions-2-routine.f90: Adjust.
* gfortran.dg/goacc/orphan-reductions-1.f90: Adjust.
* gfortran.dg/goacc/orphan-reductions-2.f90: Adjust.
---
gcc/omp-offload.c | 7 ------
.../goacc/nested-reductions-1-routine.c | 10 +++++---
.../goacc/nested-reductions-2-routine.c | 17 +++++++------
.../c-c++-common/goacc/orphan-reductions-2.c | 24 +++++++++++--------
.../goacc/nested-reductions-1-routine.f90 | 10 +++++---
.../goacc/nested-reductions-2-routine.f90 | 17 +++++++------
.../gfortran.dg/goacc/orphan-reductions-1.f90 | 4 ++++
.../gfortran.dg/goacc/orphan-reductions-2.f90 | 24 +++++++++++--------
8 files changed, 62 insertions(+), 51 deletions(-)
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 5110a424584..5cdb57d9132 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1623,13 +1623,6 @@ oacc_loop_auto_partitions (oacc_loop *loop, unsigned outer_mask,
non-innermost available level. */
unsigned this_mask = GOMP_DIM_MASK (GOMP_DIM_GANG);
- /* Orphan reductions cannot have gang partitioning. */
- if ((loop->flags & OLF_REDUCTION)
- && oacc_get_fn_attrib (current_function_decl)
- && !lookup_attribute ("omp target entrypoint",
- DECL_ATTRIBUTES (current_function_decl)))
- this_mask = GOMP_DIM_MASK (GOMP_DIM_WORKER);
-
/* Find the first outermost available partition. */
while (this_mask <= outer_mask)
this_mask <<= 1;
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-routine.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-routine.c
index 9e34614eb15..45b8cf3451f 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-routine.c
@@ -8,18 +8,21 @@ void acc_routine (void)
int i, j, k, sum, diff;
{
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
for (j = 0; j < 10; j++)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop collapse(2) reduction(+:sum)
for (i = 0; i < 10; i++)
for (j = 0; j < 10; j++)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(+:sum)
@@ -27,6 +30,7 @@ void acc_routine (void)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop collapse(2) reduction(+:sum)
@@ -34,6 +38,7 @@ void acc_routine (void)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
for (j = 0; j < 10; j++)
@@ -41,27 +46,26 @@ void acc_routine (void)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(+:sum)
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum) reduction(-:diff)
for (i = 0; i < 10; i++)
{
#pragma acc loop reduction(+:sum)
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(-:diff)
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(-:diff)
for (k = 0; k < 10; k++)
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c
index 9bd79dea4cf..3b2b0275ec8 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c
@@ -8,29 +8,29 @@ void acc_routine (void)
int i, j, k, l, sum, diff;
{
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
for (k = 0; k < 10; k++)
#pragma acc loop reduction(+:sum)
for (l = 0; l < 10; l++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
// { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -39,28 +39,28 @@ void acc_routine (void)
for (l = 0; l < 10; l++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(-:sum)
for (k = 0; k < 10; k++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
// { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -69,10 +69,10 @@ void acc_routine (void)
for (l = 0; l < 10; l++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." })
// { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -81,18 +81,17 @@ void acc_routine (void)
for (l = 0; l < 10; l++)
sum = 1;
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
#pragma acc loop reduction(+:sum) reduction(-:diff)
for (i = 0; i < 10; i++)
{
#pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
- // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(-:diff)
for (k = 0; k < 10; k++)
diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
index 941e5c6126a..6ff8698a35c 100644
--- a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
@@ -1,17 +1,20 @@
-/* Ensure that the middle end does not assign gang level parallelism
- to orphan loop containing reductions. */
+/* Verify that we diagnose "gang reduction on an orphan loop" for automatically
+ assigned gang level of parallelism. */
/* { dg-do compile } */
/* { dg-additional-options "-fopt-info-optimized-omp" } */
/* { dg-additional-options "-Wopenacc-parallelism" } */
#pragma acc routine gang
+/* { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .+3 }
+ TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it. */
int
-f1 () /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */
+f1 ()
{
int sum = 0, i;
-#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC worker vector loop parallelism" } */
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC gang vector loop parallelism" } */
for (i = 0; i < 100; i++)
sum++;
@@ -20,11 +23,12 @@ f1 () /* { dg-warning "region is gang partitioned but does not contain gang part
#pragma acc routine gang
int
-f2 () /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */
+f2 ()
{
int sum = 0, i, j;
-#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC worker loop parallelism" } */
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC gang worker loop parallelism" } */
for (i = 0; i < 100; i++)
#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC vector loop parallelism" } */
for (j = 0; j < 100; j++)
@@ -35,14 +39,14 @@ f2 () /* { dg-warning "region is gang partitioned but does not contain gang part
#pragma acc routine gang
int
-f3 () /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */
+f3 ()
{
int sum = 0, i, j, k;
-#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC worker loop parallelism" } */
+ /* { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 } */
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC gang loop parallelism" } */
for (i = 0; i < 100; i++)
-#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC seq loop parallelism" } */
- /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC worker loop parallelism" } */
for (j = 0; j < 100; j++)
#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC vector loop parallelism" } */
for (k = 0; k < 100; k++)
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90
index e8264114714..e1b0a0202d4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90
@@ -8,6 +8,7 @@ subroutine acc_routine ()
integer :: i, j, k, sum, diff
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
do j = 1, 10
@@ -17,6 +18,7 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop collapse(2) reduction(+:sum)
do i = 1, 10
do j = 1, 10
@@ -26,6 +28,7 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
@@ -36,6 +39,7 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop collapse(2) reduction(+:sum)
@@ -46,6 +50,7 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
do j = 1, 10
@@ -56,10 +61,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -68,10 +73,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum) reduction(-:diff)
do i = 1, 10
!$acc loop reduction(+:sum)
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -80,7 +85,6 @@ subroutine acc_routine ()
end do
!$acc loop reduction(-:diff)
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(-:diff)
do k = 1, 10
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90
index 98b1aa641c0..73a05207910 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90
@@ -7,10 +7,10 @@ subroutine acc_routine ()
!$acc routine gang
integer :: i, j, k, l, sum, diff
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -19,10 +19,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
do k = 1, 10
!$acc loop reduction(+:sum)
@@ -33,10 +33,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -49,10 +49,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do k = 1, 10
@@ -61,10 +61,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(-:sum)
do k = 1, 10
@@ -73,10 +73,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -89,10 +89,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -105,10 +105,10 @@ subroutine acc_routine ()
end do
end do
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
!$acc loop reduction(+:sum) reduction(-:diff)
do i = 1, 10
!$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -117,7 +117,6 @@ subroutine acc_routine ()
end do
!$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(-:diff)
do k = 1, 10
diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90
index 464dee1260a..8eed080a128 100644
--- a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90
@@ -42,6 +42,7 @@ subroutine s2
end do
!$acc loop reduction(+:sum)
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do i = 1, n
!$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" }
do j = 1, n
@@ -92,6 +93,7 @@ integer function f2 ()
end do
!$acc loop reduction(+:sum)
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do i = 1, n
!$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" }
do j = 1, n
@@ -144,6 +146,7 @@ contains
end do
!$acc loop reduction(+:sum)
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do i = 1, n
!$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" }
do j = 1, n
@@ -194,6 +197,7 @@ contains
end do
!$acc loop reduction(+:sum)
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do i = 1, n
!$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" }
do j = 1, n
diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
index 7ff0a57e620..47f8a326607 100644
--- a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
@@ -1,29 +1,33 @@
-! Ensure that the middle end does not assign gang level parallelism to
-! orphan loop containing reductions.
+! Verify that we diagnose "gang reduction on an orphan loop" for automatically
+! assigned gang level of parallelism.
! { dg-do compile }
! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-Wopenacc-parallelism" }
-subroutine s1 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" }
+subroutine s1
implicit none
!$acc routine gang
+ ! { dg-bogus "\[Ww\]arning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .-3 }
+ !TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it.
integer i, sum
sum = 0
- !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker vector loop parallelism" }
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC gang vector loop parallelism" }
do i = 1, 10
sum = sum + 1
end do
end subroutine s1
-subroutine s2 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" }
+subroutine s2
implicit none
!$acc routine gang
integer i, j, sum
sum = 0
- !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker loop parallelism" }
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC gang worker loop parallelism" }
do i = 1, 10
!$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC vector loop parallelism" }
do j = 1, 10
@@ -32,16 +36,16 @@ subroutine s2 ! { dg-warning "region is gang partitioned but does not contain ga
end do
end subroutine s2
-subroutine s3 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" }
+subroutine s3
implicit none
!$acc routine gang
integer i, j, k, sum
sum = 0
- !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker loop parallelism" }
+ ! { dg-error "gang reduction on an orphan loop" "" { target *-*-* } .+1 }
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC gang loop parallelism" }
do i = 1, 10
- !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC seq loop parallelism" }
- ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker loop parallelism" }
do j = 1, 10
!$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC vector loop parallelism" }
do k = 1, 10
--
2.33.0
