From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id EA8C0385AC26; Tue, 30 Nov 2021 12:00:17 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org EA8C0385AC26 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-5620] Re OpenACC "gang reduction on an orphan loop" error message X-Act-Checkin: gcc X-Git-Author: Frederik Harwath X-Git-Refname: refs/heads/master X-Git-Oldrev: f1a58ab0db20c0862e8b5039bd448fc8c9799cac X-Git-Newrev: c4f4c60457d1657cbd72015de3d818eb6462a0e9 Message-Id: <20211130120017.EA8C0385AC26@sourceware.org> Date: Tue, 30 Nov 2021 12:00:17 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 30 Nov 2021 12:00:18 -0000 https://gcc.gnu.org/g:c4f4c60457d1657cbd72015de3d818eb6462a0e9 commit r12-5620-gc4f4c60457d1657cbd72015de3d818eb6462a0e9 Author: Frederik Harwath Date: Mon Jul 20 11:24:21 2020 +0200 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 Diff: --- 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(-) 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