From e8c45244e5f2b69cb3b4f867b8badcedb0730987 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Tue, 2 May 2017 01:20:42 +0000 Subject: Make OpenACC orphan gang reductions errors. 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. git-svn-id: https://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@247461 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 4 + gcc/c/ChangeLog.gomp | 5 + gcc/c/c-typeck.c | 8 + gcc/cp/ChangeLog.gomp | 5 + gcc/cp/semantics.c | 8 + gcc/fortran/ChangeLog.gomp | 2 + gcc/fortran/openmp.c | 12 ++ gcc/omp-low.c | 16 +- gcc/testsuite/ChangeLog.gomp | 6 + .../c-c++-common/goacc/orphan-reductions-1.c | 58 ++++++ .../c-c++-common/goacc/orphan-reductions-2.c | 87 +++++++++ gcc/testsuite/c-c++-common/goacc/routine-4.c | 8 +- gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 2 +- .../gfortran.dg/goacc/orphan-reductions-1.f90 | 206 +++++++++++++++++++++ .../gfortran.dg/goacc/orphan-reductions-2.f90 | 85 +++++++++ 15 files changed, 505 insertions(+), 7 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 diff --git a/gcc/ChangeLog.gomp b/gcc/ChangeLog.gomp index 5aa77b0bd08..41b7454648f 100644 --- a/gcc/ChangeLog.gomp +++ b/gcc/ChangeLog.gomp @@ -3,6 +3,10 @@ * omp-low.c (verify_oacc_routine_clauses): Emit a warning when the user doesn't supply a gang, worker, vector or seq clause to an OpenACC routine construct. + (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. 2017-04-24 Cesar Philippidis diff --git a/gcc/c/ChangeLog.gomp b/gcc/c/ChangeLog.gomp index d8f2257f493..9bf1d55d9c9 100644 --- a/gcc/c/ChangeLog.gomp +++ b/gcc/c/ChangeLog.gomp @@ -1,3 +1,8 @@ +2017-05-01 Cesar Philippidis + + * c-typeck.c (c_finish_omp_clauses): Emit an error on orphan OpenACC + gang reductions. + 2017-04-26 Cesar Philippidis * c-parser.c (c_parser_omp_variable_list): New c_omp_region_type diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 61a95b06372..b04db4449df 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_region_type ort) goto check_dup_generic; case OMP_CLAUSE_REDUCTION: + if (ort == 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 = true; + break; + } need_implicitly_determined = true; t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) diff --git a/gcc/cp/ChangeLog.gomp b/gcc/cp/ChangeLog.gomp index 1bab24d1e83..431aeb33eed 100644 --- a/gcc/cp/ChangeLog.gomp +++ b/gcc/cp/ChangeLog.gomp @@ -1,3 +1,8 @@ +2017-05-01 Cesar Philippidis + + * semantics.c (finish_omp_clauses): Emit an error on orphan OpenACC + gang reductions. + 2017-04-26 Cesar Philippidis * parser.c (cp_parser_omp_var_list_no_open): New c_omp_region_type diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 9760f07e68d..6e8fb171fd1 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 = ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP); goto check_dup_generic; case OMP_CLAUSE_REDUCTION: + if (ort == 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 = true; + break; + } field_ok = ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP); t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) diff --git a/gcc/fortran/ChangeLog.gomp b/gcc/fortran/ChangeLog.gomp index 031e10b08c6..2f188150a28 100644 --- a/gcc/fortran/ChangeLog.gomp +++ b/gcc/fortran/ChangeLog.gomp @@ -6,6 +6,8 @@ (gfc_match_oacc_routine): Emit a warning when the user doesn't supply a gang, worker, vector or seq clause to an OpenACC routine construct. + (resolve_oacc_loop_blocks): Emit an error on orphan OpenACC + gang reductions. 2017-04-27 Cesar Philippidis diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 88ccff2cd0a..77a9c8d0eb6 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -6099,6 +6099,18 @@ resolve_oacc_loop_blocks (gfc_code *code) break; } + if (code->op == EXEC_OACC_LOOP + && code->ext.omp_clauses->lists[OMP_LIST_REDUCTION] + && code->ext.omp_clauses->gang) + { + for (c = omp_current_ctx; c; c = c->previous) + if (!oacc_is_loop (c->code)) + break; + if (c == 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 88a1bb9c5d7..cf299c12ecc 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -272,9 +272,10 @@ enum oacc_loop_flags { OLF_INDEPENDENT = 1u << 2, /* Iterations are known independent. */ OLF_GANG_STATIC = 1u << 3, /* Gang partitioning is static (has op). */ OLF_TILE = 1u << 4, /* Tiled loop. */ + OLF_REDUCTION = 1u << 5, /* Reduction loop. */ /* Explicitly specified loop axes. */ - OLF_DIM_BASE = 5, + OLF_DIM_BASE = 6, OLF_DIM_GANG = 1u << (OLF_DIM_BASE + GOMP_DIM_GANG), OLF_DIM_WORKER = 1u << (OLF_DIM_BASE + GOMP_DIM_WORKER), OLF_DIM_VECTOR = 1u << (OLF_DIM_BASE + GOMP_DIM_VECTOR), @@ -6616,6 +6617,10 @@ lower_oacc_head_mark (location_t loc, tree ddvar, tree clauses, tag |= OLF_TILE; break; + case OMP_CLAUSE_REDUCTION: + tag |= OLF_REDUCTION; + break; + case OMP_CLAUSE_DEVICE_TYPE: /* TODO: Add device type handling. */ goto done; @@ -20944,7 +20949,14 @@ oacc_loop_auto_partitions (oacc_loop *loop, unsigned outer_mask, /* Allocate outermost and non-innermost loops at the outermost non-innermost available level. */ unsigned this_mask = GOMP_DIM_MASK (GOMP_DIM_GANG); - + + /* 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 = 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/ChangeLog.gomp b/gcc/testsuite/ChangeLog.gomp index fbcdf4ece83..80876c0b569 100644 --- a/gcc/testsuite/ChangeLog.gomp +++ b/gcc/testsuite/ChangeLog.gomp @@ -22,6 +22,12 @@ * gfortran.dg/goacc/routine-9.f90: Likewise. * gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise. * gfortran.dg/goacc/routine-without-clauses.f90: New test. + * 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. 2017-04-27 Cesar Philippidis diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c new file mode 100644 index 00000000000..2a5825e9599 --- /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 = 0; +#pragma acc loop seq reduction(+:sum) + for (i = 0; i < n; i++) + sum = sum + 1; + + return sum; +} + +#pragma acc routine gang +int +gang_reduction (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc loop gang reduction(+:s1) /* { dg-error "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-error "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + + + return s1 + s2; +} + +#pragma acc routine worker +int +worker_reduction (int n) +{ + int i, sum = 0; +#pragma acc loop worker reduction(+:sum) + for (i = 0; i < n; i++) + sum = sum + 3; + + return sum; +} + +#pragma acc routine vector +int +vector_reduction (int n) +{ + int i, sum = 0; +#pragma acc loop vector reduction(+:sum) + for (i = 0; i < n; i++) + sum = sum + 4; + + return sum; +} diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c new file mode 100644 index 00000000000..51d2596a8d2 --- /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 gang partitioned code" } */ +{ + int sum = 0, i; + +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (i = 0; i < 100; i++) + sum++; + + return sum; +} + +#pragma acc routine gang +int +f2 () /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */ +{ + int sum = 0, i, j; + +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (i = 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (j = 0; j < 100; j++) + sum++; + + return sum; +} + +#pragma acc routine gang +int +f3 () /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */ +{ + int sum = 0, i, j, k; + +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (i = 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (j = 0; j < 100; j++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (k = 0; k < 100; k++) + sum++; + + return sum; +} + +int +main () +{ + int sum = 0, i, j, k; + +#pragma acc parallel copy (sum) + { +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (i = 0; i < 100; i++) + sum++; + } + +#pragma acc parallel copy (sum) + { +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (i = 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (j = 0; j < 100; j++) + sum++; + } + +#pragma acc parallel copy (sum) + { +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (i = 0; i < 100; i++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (j = 0; j < 100; j++) +#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism " } */ + for (k = 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 3e5fc4f5d77..0bead00fcca 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 = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" } +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; @@ -48,7 +48,7 @@ void vector (void) /* { dg-message "declared here" 1 } */ for (int i = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" } +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; @@ -74,7 +74,7 @@ void worker (void) /* { dg-message "declared here" 2 } */ for (int i = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" } +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; @@ -100,7 +100,7 @@ void gang (void) /* { dg-message "declared here" 3 } */ for (int i = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c index ac886c78f95..85e73b14cee 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) } } -/* { dg-final { scan-tree-dump "OpenACC loops.*Loop 0\\\(0\\\).*Loop 24\\\(1\\\).*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 36\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 36\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Loop 6\\\(6\\\).*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*Head-1:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-1:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(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_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 68\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 68\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Loop 6\\\(6\\\).*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*Head-1:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-1:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(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 00000000000..c7fcc9d4ac5 --- /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 = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel +end subroutine s1 + +subroutine s2 + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do +end subroutine s2 + +integer function f1 () + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + f1 = sum +end function f1 + +integer function f2 () + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do + + f2 = sum +end function f2 + +module m +contains + subroutine s3 + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + end subroutine s3 + + subroutine s4 + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do + end subroutine s4 + + integer function f3 () + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + f3 = sum + end function f3 + + integer function f4 () + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do + + f4 = 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 00000000000..8ec60cbe55f --- /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 contain gang partitioned code" } + implicit none + !$acc routine gang + integer i, sum + + !$acc loop reduction (+:sum) ! { dg-message "Detected 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" } + implicit none + !$acc routine gang + integer i, j, sum + + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i = 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j = 1, 10 + sum = sum + 1 + end do + end do +end subroutine s2 + +subroutine s3 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" } + implicit none + !$acc routine gang + integer i, j, k, sum + + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i = 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j = 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do k = 1, 10 + sum = 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 = 1, 10 + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel copy(sum) + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i = 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j = 1, 10 + sum = sum + 1 + end do + end do + !$acc end parallel + + !$acc parallel copy(sum) + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do i = 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do j = 1, 10 + !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism " } + do k = 1, 10 + sum = sum + 1 + end do + end do + end do + !$acc end parallel +end subroutine s4 + +! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } 39 } -- cgit v1.2.3