From c4f4c60457d1657cbd72015de3d818eb6462a0e9 Mon Sep 17 00:00:00 2001 From: Frederik Harwath Date: Mon, 20 Jul 2020 11:24:21 +0200 Subject: [PATCH] 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 --- gcc/fortran/openmp.c | 9 +- .../c-c++-common/goacc/orphan-reductions-3.c | 102 ++++++++++++++++++ .../gfortran.dg/goacc/orphan-reductions-3.f90 | 89 +++++++++++++++ 3 files changed, 196 insertions(+), 4 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-3.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-3.f90 diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index b4100577e51..7950c7fb43d 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -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); } diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-3.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-3.c new file mode 100644 index 00000000000..cd8ad274ebb --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-3.c @@ -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; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-3.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-3.f90 new file mode 100644 index 00000000000..1e0b1d64578 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-3.f90 @@ -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 -- 2.33.0