aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorCesar Philippidis <cesar@codesourcery.com>2017-05-02 01:20:42 +0000
committerCesar Philippidis <cesar@codesourcery.com>2017-05-02 01:20:42 +0000
commite8c45244e5f2b69cb3b4f867b8badcedb0730987 (patch)
treea3ec66b342296c22b74fc022d164c26c219cc409
parent4baac306c48ad8d44b3b827942d3a444150aaf73 (diff)
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
-rw-r--r--gcc/ChangeLog.gomp4
-rw-r--r--gcc/c/ChangeLog.gomp5
-rw-r--r--gcc/c/c-typeck.c8
-rw-r--r--gcc/cp/ChangeLog.gomp5
-rw-r--r--gcc/cp/semantics.c8
-rw-r--r--gcc/fortran/ChangeLog.gomp2
-rw-r--r--gcc/fortran/openmp.c12
-rw-r--r--gcc/omp-low.c16
-rw-r--r--gcc/testsuite/ChangeLog.gomp6
-rw-r--r--gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c58
-rw-r--r--gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c87
-rw-r--r--gcc/testsuite/c-c++-common/goacc/routine-4.c8
-rw-r--r--gcc/testsuite/gcc.dg/goacc/loop-processing-1.c2
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90206
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f9085
15 files changed, 505 insertions, 7 deletions
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 <cesar@codesourcery.com>
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 <cesar@codesourcery.com>
+
+ * c-typeck.c (c_finish_omp_clauses): Emit an error on orphan OpenACC
+ gang reductions.
+
2017-04-26 Cesar Philippidis <cesar@codesourcery.com>
* 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 <cesar@codesourcery.com>
+
+ * semantics.c (finish_omp_clauses): Emit an error on orphan OpenACC
+ gang reductions.
+
2017-04-26 Cesar Philippidis <cesar@codesourcery.com>
* 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 <cesar@codesourcery.com>
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 <cesar@codesourcery.com>
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 <assert.h>
+
+#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 <acc loop worker vector>" } */
+ 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 <acc loop worker>" } */
+ for (i = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */
+ 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 <acc loop worker>" } */
+ for (i = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop seq>" } */
+ for (j = 0; j < 100; j++)
+#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */
+ 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 <acc loop gang vector>" } */
+ for (i = 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 = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */
+ for (j = 0; j < 100; j++)
+ sum++;
+ }
+
+#pragma acc parallel copy (sum)
+ {
+#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop gang>" } */
+ for (i = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop worker>" } */
+ for (j = 0; j < 100; j++)
+#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */
+ 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 <acc loop worker vector>" }
+ 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 <acc loop worker>" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" }
+ 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 <acc loop worker>" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop seq>" }
+ do j = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" }
+ 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 <acc loop gang vector>" }
+ do i = 1, 10
+ sum = sum + 1
+ end do
+ !$acc end parallel
+
+ !$acc parallel copy(sum)
+ !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop gang worker>" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" }
+ 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 <acc loop gang>" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop worker>" }
+ do j = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" }
+ 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 }