From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 37134 invoked by alias); 2 May 2017 01:32:16 -0000 Mailing-List: contact fortran-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Subscribe: List-Post: List-Help: , Sender: fortran-owner@gcc.gnu.org Received: (qmail 37110 invoked by uid 89); 2 May 2017 01:32:16 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.7 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy=gang, Detected, accord, insufficient X-Spam-User: qpsmtpd, 2 recipients X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 02 May 2017 01:32:13 +0000 Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204]) by relay1.mentorg.com with esmtp id 1d5Mfk-000501-1E from Cesar_Philippidis@mentor.com ; Mon, 01 May 2017 18:32:12 -0700 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Mon, 1 May 2017 18:28:00 -0700 From: Cesar Philippidis Subject: [gomp4] Make OpenACC orphan gang reductions errors To: "gcc-patches@gcc.gnu.org" , Fortran List Message-ID: <91521931-373f-a3a6-7a3f-1752f08858a1@codesourcery.com> Date: Tue, 02 May 2017 01:32:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.8.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------08496B6434AFAC5741201C69" X-ClientProxiedBy: svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) To SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) X-IsSubscribed: yes X-SW-Source: 2017-05/txt/msg00004.txt.bz2 --------------08496B6434AFAC5741201C69 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 571 This patch promotes all OpenACC gang reductions on orphan loops as errors. Accord to the spec, orphan loops are those which are not lexically nested inside an OpenACC parallel or kernels regions. I.e., acc loops inside acc routines. At first I thought this could be a warning because the gang reduction finalizer uses an atomic update. However, because there is no synchronization between gangs, there is way to guarantee that reduction will have completed once a single gang entity returns from the acc routine call. I've applied this patch to gomp-4_0-branch. Cesar --------------08496B6434AFAC5741201C69 Content-Type: text/x-patch; name="gomp4-orphan-reductions.diff" Content-Transfer-Encoding: quoted-printable Content-Disposition: attachment; filename="gomp4-orphan-reductions.diff" Content-length: 21033 2017-05-01 Cesar Philippidis gcc/c/ * c-typeck.c (c_finish_omp_clauses): Emit an error on orphan OpenACC gang reductions. gcc/cp/ * semantics.c (finish_omp_clauses): Emit an error on orphan OpenACC gang reductions. gcc/fortran/ * openmp.c (resolve_oacc_loop_blocks): Emit an error on orphan OpenACC gang reductions. gcc/ * omp-low.c (enum oacc_loop_flags): Add OLF_REDUCTION enum. (lower_oacc_head_mark): Use it to mark OpenACC reductions. (oacc_loop_auto_partitions): Don't assign gang level parallelism to orphan reductions. gcc/testsuite/ * c-c++-common/goacc/orphan-reductions-1.c: New test. * c-c++-common/goacc/orphan-reductions-2.c: New test. * c-c++-common/goacc/routine-4.c: Update test case. * gcc.dg/goacc/loop-processing-1.c: Likewise. * gfortran.dg/goacc/orphan-reductions-1.f90: New test. * gfortran.dg/goacc/orphan-reductions-2.f90: New test. diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 61a95b0..b04db44 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12602,6 +12602,14 @@ c_finish_omp_clauses (tree clauses, enum c_omp_reg= ion_type ort) goto check_dup_generic; =20 case OMP_CLAUSE_REDUCTION: + if (ort =3D=3D C_ORT_ACC && get_oacc_fn_attrib (current_function_decl) + && find_omp_clause (clauses, OMP_CLAUSE_GANG)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "gang reduction on an orphan loop"); + remove =3D true; + break; + } need_implicitly_determined =3D true; t =3D OMP_CLAUSE_DECL (c); if (TREE_CODE (t) =3D=3D TREE_LIST) diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 9760f07..6e8fb17 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5870,6 +5870,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_= type ort) field_ok =3D ((ort & C_ORT_OMP_DECLARE_SIMD) =3D=3D C_ORT_OMP); goto check_dup_generic; case OMP_CLAUSE_REDUCTION: + if (ort =3D=3D C_ORT_ACC && get_oacc_fn_attrib (current_function_decl) + && find_omp_clause (clauses, OMP_CLAUSE_GANG)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "gang reduction on an orphan loop"); + remove =3D true; + break; + } field_ok =3D ((ort & C_ORT_OMP_DECLARE_SIMD) =3D=3D C_ORT_OMP); t =3D OMP_CLAUSE_DECL (c); if (TREE_CODE (t) =3D=3D TREE_LIST) diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 72c6669..fb51b40 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -6090,6 +6090,18 @@ resolve_oacc_loop_blocks (gfc_code *code) break; } =20 + if (code->op =3D=3D EXEC_OACC_LOOP + && code->ext.omp_clauses->lists[OMP_LIST_REDUCTION] + && code->ext.omp_clauses->gang) + { + for (c =3D omp_current_ctx; c; c =3D c->previous) + if (!oacc_is_loop (c->code)) + break; + if (c =3D=3D NULL || !(oacc_is_parallel (c->code) + || oacc_is_kernels (c->code))) + gfc_error ("gang reduction on an orphan loop at %L", &code->loc); + } + if (code->ext.omp_clauses->seq) { if (code->ext.omp_clauses->independent) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index cc209ba..d6c62f9 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -272,9 +272,10 @@ enum oacc_loop_flags { OLF_INDEPENDENT =3D 1u << 2, /* Iterations are known independent. */ OLF_GANG_STATIC =3D 1u << 3, /* Gang partitioning is static (has op). */ OLF_TILE =3D 1u << 4, /* Tiled loop. */ + OLF_REDUCTION =3D 1u << 5, /* Reduction loop. */ =20=20=20 /* Explicitly specified loop axes. */ - OLF_DIM_BASE =3D 5, + OLF_DIM_BASE =3D 6, OLF_DIM_GANG =3D 1u << (OLF_DIM_BASE + GOMP_DIM_GANG), OLF_DIM_WORKER =3D 1u << (OLF_DIM_BASE + GOMP_DIM_WORKER), OLF_DIM_VECTOR =3D 1u << (OLF_DIM_BASE + GOMP_DIM_VECTOR), @@ -6616,6 +6617,10 @@ lower_oacc_head_mark (location_t loc, tree ddvar, tr= ee clauses, tag |=3D OLF_TILE; break; =20 + case OMP_CLAUSE_REDUCTION: + tag |=3D OLF_REDUCTION; + break; + case OMP_CLAUSE_DEVICE_TYPE: /* TODO: Add device type handling. */ goto done; @@ -20942,7 +20947,14 @@ oacc_loop_auto_partitions (oacc_loop *loop, unsign= ed outer_mask, /* Allocate outermost and non-innermost loops at the outermost non-innermost available level. */ unsigned this_mask =3D GOMP_DIM_MASK (GOMP_DIM_GANG); -=20=20=20=20=20=20 + + /* Orphan reductions cannot have gang partitioning. */ + if ((loop->flags & OLF_REDUCTION) + && get_oacc_fn_attrib (current_function_decl) + && !lookup_attribute ("omp target entrypoint", + DECL_ATTRIBUTES (current_function_decl))) + this_mask =3D GOMP_DIM_MASK (GOMP_DIM_WORKER); + /* Find the first outermost available partition. */ while (this_mask <=3D outer_mask) this_mask <<=3D 1; diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c b/gcc/t= estsuite/c-c++-common/goacc/orphan-reductions-1.c new file mode 100644 index 0000000..2a5825e --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c @@ -0,0 +1,58 @@ +/* Test orphan reductions. */ + +/* { dg-do compile } */ + +#include + +#pragma acc routine seq +int +seq_reduction (int n) +{ + int i, sum =3D 0; +#pragma acc loop seq reduction(+:sum) + for (i =3D 0; i < n; i++) + sum =3D sum + 1; + + return sum; +} + +#pragma acc routine gang +int +gang_reduction (int n) +{ + int i, s1 =3D 0, s2 =3D 0; +#pragma acc loop gang reduction(+:s1) /* { dg-error "gang reduction on an = orphan loop" } */ + for (i =3D 0; i < n; i++) + s1 =3D s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-error "gang reduction on an = orphan loop" } */ + for (i =3D 0; i < n; i++) + s2 =3D s2 + 2; + + + return s1 + s2; +} + +#pragma acc routine worker +int +worker_reduction (int n) +{ + int i, sum =3D 0; +#pragma acc loop worker reduction(+:sum) + for (i =3D 0; i < n; i++) + sum =3D sum + 3; + + return sum; +} + +#pragma acc routine vector +int +vector_reduction (int n) +{ + int i, sum =3D 0; +#pragma acc loop vector reduction(+:sum) + for (i =3D 0; i < n; i++) + sum =3D sum + 4; + + return sum; +} diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c b/gcc/t= estsuite/c-c++-common/goacc/orphan-reductions-2.c new file mode 100644 index 0000000..51d2596 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c @@ -0,0 +1,87 @@ +/* Ensure that the middle end does not assign gang level parallelism + to orphan loop containing reductions. */ + +/* { dg-do compile } */ +/* { dg-additional-options "-fopt-info-note-omp" } */ + +#pragma acc routine gang +int +f1 () /* { dg-warning "region is gang partitioned but does not contain gan= g partitioned code" } */ +{ + int sum =3D 0, i; + +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop worker vector>" } */ + for (i =3D 0; i < 100; i++) + sum++; + + return sum; +} + +#pragma acc routine gang +int +f2 () /* { dg-warning "region is gang partitioned but does not contain gan= g partitioned code" } */ +{ + int sum =3D 0, i, j; + +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop worker>" } */ + for (i =3D 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop vector>" } */ + for (j =3D 0; j < 100; j++) + sum++; + + return sum; +} + +#pragma acc routine gang +int +f3 () /* { dg-warning "region is gang partitioned but does not contain gan= g partitioned code" } */ +{ + int sum =3D 0, i, j, k; + +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop worker>" } */ + for (i =3D 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop seq>" } */ + for (j =3D 0; j < 100; j++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop vector>" } */ + for (k =3D 0; k < 100; k++) + sum++; + + return sum; +} + +int +main () +{ + int sum =3D 0, i, j, k; + +#pragma acc parallel copy (sum) + { +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop gang vector>" } */ + for (i =3D 0; i < 100; i++) + sum++; + } + +#pragma acc parallel copy (sum) + { +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop gang worker>" } */ + for (i =3D 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop vector>" } */ + for (j =3D 0; j < 100; j++) + sum++; + } + +#pragma acc parallel copy (sum) + { +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop gang>" } */ + for (i =3D 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop worker>" } */ + for (j =3D 0; j < 100; j++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <= acc loop vector>" } */ + for (k =3D 0; k < 100; k++) + sum++; + } + + return sum; +} + +/* { dg-warning "insufficient partitioning available to parallelize loop" = "" { target *-*-* } 43 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-4.c b/gcc/testsuite/c= -c++-common/goacc/routine-4.c index 3e5fc4f..0bead00 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-4.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-4.c @@ -22,7 +22,7 @@ void seq (void) for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by conta= ining routine" } +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 @@ -48,7 +48,7 @@ void vector (void) /* { dg-message "declared here" 1 } */ for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by conta= ining routine" } +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 @@ -74,7 +74,7 @@ void worker (void) /* { dg-message "declared here" 2 } */ for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by conta= ining routine" } +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 @@ -100,7 +100,7 @@ void gang (void) /* { dg-message "declared here" 3 } */ for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite= /gcc.dg/goacc/loop-processing-1.c index ac886c7..85e73b1 100644 --- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c +++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c @@ -15,4 +15,4 @@ void vector_1 (int *ary, int size) } } =20 -/* { dg-final { scan-tree-dump "OpenACC loops.*Loop 0\\\(0\\\).*Loop 24\\\= (1\\\).*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 36\\= \);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 1= , 36\\\);.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_FORK, \\\.data_dep= \\\.\[0-9_\]+, 0\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OA= CC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = =3D UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Loop 6\\\(6\\\= ).*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*H= ead-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\= );.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0= -9_\]+, 1\\\);.*Head-1:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD= _MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ =3D UNI= QUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-1:.*\\\.data_dep= \\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 2\\= \);.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[= 0-9_\]+, 2\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_TAI= L_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ =3D UN= IQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 1\\\);" "oaccdevlow" } } */ +/* { dg-final { scan-tree-dump "OpenACC loops.*Loop 0\\\(0\\\).*Loop 44\\\= (1\\\).*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 68\\= \);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 1= , 68\\\);.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_FORK, \\\.data_dep= \\\.\[0-9_\]+, 0\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OA= CC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = =3D UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Loop 6\\\(6\\\= ).*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*H= ead-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\= );.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0= -9_\]+, 1\\\);.*Head-1:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_HEAD= _MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ =3D UNI= QUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-1:.*\\\.data_dep= \\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 2\\= \);.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[= 0-9_\]+, 2\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ =3D UNIQUE \\\(OACC_TAI= L_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ =3D UN= IQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 1\\\);" "oaccdevlow" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 b/gcc/= testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 new file mode 100644 index 0000000..c7fcc9d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 @@ -0,0 +1,206 @@ +! Verify that gang reduction on orphan OpenACC loops reported as errors. + +! { dg-do compile } + +subroutine s1 + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel +end subroutine s1 + +subroutine s2 + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orph= an loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an o= rphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do +end subroutine s2 + +integer function f1 () + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + f1 =3D sum +end function f1 + +integer function f2 () + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orph= an loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an o= rphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do + + f2 =3D sum +end function f2 + +module m +contains + subroutine s3 + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + end subroutine s3 + + subroutine s4 + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an or= phan loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an= orphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do + end subroutine s4 + + integer function f3 () + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + f3 =3D sum + end function f3 + + integer function f4 () + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an or= phan loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an= orphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do + + f4 =3D sum + end function f4 +end module m 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 0000000..8ec60cb --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 @@ -0,0 +1,85 @@ +! Ensure that the middle end does not assign gang level parallelism to +! orphan loop containing reductions. + +! { dg-do compile } +! { dg-additional-options "-fopt-info-note-omp" } + +subroutine s1 ! { dg-warning "region is gang partitioned but does not cont= ain gang partitioned code" } + implicit none + !$acc routine gang + integer i, sum + + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i =3D 1, 10 + sum =3D sum + 1 + end do +end subroutine s1 + +subroutine s2 ! { dg-warning "region is gang partitioned but does not cont= ain gang partitioned code" } + implicit none + !$acc routine gang + integer i, j, sum + + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i =3D 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j =3D 1, 10 + sum =3D sum + 1 + end do + end do +end subroutine s2 + +subroutine s3 ! { dg-warning "region is gang partitioned but does not cont= ain gang partitioned code" } + implicit none + !$acc routine gang + integer i, j, k, sum + + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i =3D 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j =3D 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism = " } + do k =3D 1, 10 + sum =3D sum + 1 + end do + end do + end do +end subroutine s3 + +subroutine s4 + implicit none + + integer i, j, k, sum + + !$acc parallel copy(sum) + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i =3D 1, 10 + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel copy(sum) + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i =3D 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j =3D 1, 10 + sum =3D sum + 1 + end do + end do + !$acc end parallel + + !$acc parallel copy(sum) + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i =3D 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j =3D 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism = " } + do k =3D 1, 10 + sum =3D sum + 1 + end do + end do + end do + !$acc end parallel +end subroutine s4 + +! { dg-warning "insufficient partitioning available to parallelize loop" "= " { target *-*-* } 39 } --------------08496B6434AFAC5741201C69--