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>
This commit is contained in:
Frederik Harwath 2020-07-20 11:24:21 +02:00 committed by Thomas Schwinge
parent f1a58ab0db
commit c4f4c60457
3 changed files with 196 additions and 4 deletions

View File

@ -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);
}

View File

@ -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;
}

View File

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