* [gomp4] Make OpenACC orphan gang reductions errors
@ 2017-05-02 1:32 Cesar Philippidis
[not found] ` <xu8f8sfv4shz.fsf@harwath.name>
` (3 more replies)
0 siblings, 4 replies; 8+ messages in thread
From: Cesar Philippidis @ 2017-05-02 1:32 UTC (permalink / raw)
To: gcc-patches, Fortran List
[-- Attachment #1: Type: text/plain, Size: 571 bytes --]
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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: gomp4-orphan-reductions.diff --]
[-- Type: text/x-patch; name="gomp4-orphan-reductions.diff", Size: 21078 bytes --]
2017-05-01 Cesar Philippidis <cesar@codesourcery.com>
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_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/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 = ((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/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;
}
+ 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 cc209ba..d6c62f9 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;
@@ -20942,7 +20947,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/c-c++-common/goacc/orphan-reductions-1.c b/gcc/testsuite/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 <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 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 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 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 = 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 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)
}
}
-/* { 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 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 = 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 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 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 }
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message
[not found] ` <xu8f8sfv4shz.fsf@harwath.name>
@ 2020-07-07 10:42 ` Thomas Schwinge
2020-07-07 12:39 ` Frederik Harwath
2020-07-20 10:26 ` Frederik Harwath
0 siblings, 2 replies; 8+ messages in thread
From: Thomas Schwinge @ 2020-07-07 10:42 UTC (permalink / raw)
To: Frederik Harwath, Julian Brown, Kwok Cheung Yeung
Cc: fortran, gcc-patches, Tobias Burnus, Jakub Jelinek
Hi Frederik!
(CC <fortran@gcc.gnu.org> added, for everything touching gfortran.)
On 2020-07-07T10:52:08+0200, Frederik Harwath <frederik@codesourcery.com> wrote:
> This patch fixes the check for reductions on orphaned gang loops
This is the "Make OpenACC orphan gang reductions errors" functionality
originally added in gomp-4_0-branch r247461.
> the Fortran frontend which (in contrast to the C, C++ frontends)
> erroneously rejects reductions on gang loops that are contained in
> "kernels" constructs and which hence are not orphaned.
>
> According to the OpenACC standard version 2.5 and later, reductions on
> orphaned gang loops are explicitly disallowed (cf. section "Changes
> from Version 2.0 to 2.5"). Remember that a loop is "orphaned" if it is
> not lexically contained in a compute construct (cf. section "Loop
> construct" of the OpenACC standard), i.e. in either a "parallel", a
> "serial", or a "kernels" construct.
Or the other way round: a 'loop' construct is orphaned if it appears
inside a 'routine' region, right?
> The patch has been tested by running the GCC and libgomp testsuites.
> The latter tests ran with offloading to nvptx although that should not
> be important here unless there was some very subtle reason for
> forbidding the gang reductions on kernels loops. As expect, there seems
> to be no such reason, i.e. I observed no regressions with the patch.
Note that the aforementioned gomp-4_0-branch r247461,
openacc-gcc-7-branch commit 0554f9f79325960c72166327d442a553cd35bad9, and
openacc-gcc-8-branch commit 65dd9cf3b3c45d64d72967df1e4a54778cb4e35f
still do contain the appropriate 'kernels' handling. Just in
openacc-gcc-9-branch commit 533beb2ec19f8486e4b1b645a153746f96b41f04 this
got (a) mixed together with a bunch of other, unrelated changes ("Various
OpenACC reduction enhancements"), and (b) the 'kernels' handling got
removed. Julian (Git author), or Kwok (Git committer), do you remember
any rationale for that? Later, this then got picked into devel/omp/gcc-9
commit 3fa4bb72dcb3b9171952a0eca5310bb8811d5ffd, and devel/omp/gcc-10
commit 6b3e1f7f05cd360bbd356b3f78511aa2ec3f40c3.
> Can I include the patch in OG10?
Unless Julian/Kwok speak up soon: OK, thanks.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
May want to remove "libgomp" from the first line of the commit log --
this commit doesn't relate to libgomp specifically.
(Ideally, we'd also test 'serial' construct in addition to 'kernels',
'parallel', but we can add that later. I anyway have a WIP patch
waiting, adding more 'serial' construct testing, for a different reason,
so I'll include it there.)
Grüße
Thomas
> From 7320635211fff3a773beb0de1914dbfcc317ab37 Mon Sep 17 00:00:00 2001
> From: Frederik Harwath <frederik@codesourcery.com>
> Date: Tue, 7 Jul 2020 10:41:21 +0200
> Subject: [PATCH] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan
> loop" error message
>
> According to the OpenACC standard version 2.5 and later, reductions on
> orphaned gang loops are explicitly disallowed (cf. section "Changes
> from Version 2.0 to 2.5"). A loop is "orphaned" if it is not
> lexically contained in a compute construct (cf. section "Loop
> construct" of the OpenACC standard), i.e. in either a "parallel", a
> "serial", or a "kernels" construct.
>
> This commit fixes the check for reductions on orphaned gang loops in
> the Fortran frontend which (in contrast to the C, C++ frontends)
> erroneously rejects reductions on gang loops that are contained in
> "kernels" constructs.
>
> 2020-07-07 Frederik Harwath <frederik@codesourcery.com>
>
> gcc/fortran/
>
> * openmp.c (oacc_is_parallel_or_serial): Removed function.
> (oacc_is_kernels): New function.
> (oacc_is_compute_construct): New 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-2.f90: New test
> verifying that the error message is not emitted for
> non-orphaned loops.
>
> * c-c++-common/goacc/orphan-reductions-2.c: Likewise for C and C++.
> ---
> gcc/fortran/openmp.c | 13 +++-
> .../c-c++-common/goacc/orphan-reductions-2.c | 69 +++++++++++++++++++
> .../gfortran.dg/goacc/orphan-reductions-2.f90 | 58 ++++++++++++++++
> 3 files changed, 137 insertions(+), 3 deletions(-)
> create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
> create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
>
> diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
> index 28408c4c99a..83c498112a8 100644
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -5926,9 +5926,16 @@ oacc_is_serial (gfc_code *code)
> }
>
> static bool
> -oacc_is_parallel_or_serial (gfc_code *code)
> +oacc_is_kernels (gfc_code *code)
> {
> - return oacc_is_parallel (code) || oacc_is_serial (code);
> + return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP;
> +}
> +
> +static bool
> +oacc_is_compute_construct (gfc_code *code)
> +{
> + return oacc_is_parallel (code) || oacc_is_serial (code)
> + || oacc_is_kernels (code);
> }
>
> static gfc_statement
> @@ -6222,7 +6229,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))
> + 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-2.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
> new file mode 100644
> index 00000000000..2b651fd2b9f
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
> @@ -0,0 +1,69 @@
> +/* Verify that the error message for gang reduction on orphaned OpenACC loops
> + is not reported for non-orphaned loops. */
> +
> +#include <assert.h>
> +
> +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
> +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-2.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
> new file mode 100644
> index 00000000000..13887a059fe
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
> @@ -0,0 +1,58 @@
> +! Verify that the error message for gang reductions on orphaned OpenACC loops
> +! is not reported for non-orphaned loops.
> +
> +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 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
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message
2020-07-07 10:42 ` [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message Thomas Schwinge
@ 2020-07-07 12:39 ` Frederik Harwath
2020-07-20 10:26 ` Frederik Harwath
1 sibling, 0 replies; 8+ messages in thread
From: Frederik Harwath @ 2020-07-07 12:39 UTC (permalink / raw)
To: Thomas Schwinge, Julian Brown, Kwok Cheung Yeung
Cc: fortran, gcc-patches, Tobias Burnus, Jakub Jelinek
Thomas Schwinge <thomas@codesourcery.com> writes:
Hi Thomas,
> (CC <fortran@gcc.gnu.org> added, for everything touching gfortran.)
Thanks!
> On 2020-07-07T10:52:08+0200, Frederik Harwath <frederik@codesourcery.com> wrote:
>> This patch fixes the check for reductions on orphaned gang loops
>
> This is the "Make OpenACC orphan gang reductions errors" functionality
> originally added in gomp-4_0-branch r247461.
>
>> the Fortran frontend which (in contrast to the C, C++ frontends)
>> erroneously rejects reductions on gang loops that are contained in
>> "kernels" constructs and which hence are not orphaned.
>>
>> According to the OpenACC standard version 2.5 and later, reductions on
>> orphaned gang loops are explicitly disallowed (cf. section "Changes
>> from Version 2.0 to 2.5"). Remember that a loop is "orphaned" if it is
>> not lexically contained in a compute construct (cf. section "Loop
>> construct" of the OpenACC standard), i.e. in either a "parallel", a
>> "serial", or a "kernels" construct.
>
> Or the other way round: a 'loop' construct is orphaned if it appears
> inside a 'routine' region, right?
The "not lexically contained in a compute construct" definition is
from the standard. Assuming that the frontend's parser rejects "loop"
directives if they do not occur inside of either the "serial",
"parallel", "kernels" compute constructs or in a function with a
"routine" directive, both definitions should be indeed equivalent ;-).
> Unless Julian/Kwok speak up soon: OK, thanks.
>
> Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
>
> May want to remove "libgomp" from the first line of the commit log --
> this commit doesn't relate to libgomp specifically.
Right.
> (Ideally, we'd also test 'serial' construct in addition to 'kernels',
> 'parallel', but we can add that later. I anyway have a WIP patch
> waiting, adding more 'serial' construct testing, for a different reason,
> so I'll include it there.)
I had left this out intentionally, because having the gang reduction in
the serial construct leads to a "region contains gang partitioned
code but is not gang partitioned"
error. Of course, we might still add a test case with that expectation.
Thanks for the review!
Frederik
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message
2020-07-07 10:42 ` [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message Thomas Schwinge
2020-07-07 12:39 ` Frederik Harwath
@ 2020-07-20 10:26 ` Frederik Harwath
2021-11-30 12:12 ` Thomas Schwinge
1 sibling, 1 reply; 8+ messages in thread
From: Frederik Harwath @ 2020-07-20 10:26 UTC (permalink / raw)
To: Thomas Schwinge, Julian Brown, Kwok Cheung Yeung
Cc: fortran, gcc-patches, Tobias Burnus, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 963 bytes --]
Thomas Schwinge <thomas@codesourcery.com> writes:
Hi Thomas,
>> Can I include the patch in OG10?
>
> Unless Julian/Kwok speak up soon: OK, thanks.
This has been delayed a bit by my vacation, but I have now committed
the patch.
> May want to remove "libgomp" from the first line of the commit log --
> this commit doesn't relate to libgomp specifically.
>
> (Ideally, we'd also test 'serial' construct in addition to 'kernels',
> 'parallel', but we can add that later. I anyway have a WIP patch
> waiting, adding more 'serial' construct testing, for a different reason,
> so I'll include it there.)
I forgot to remove "libgomp" from the commit message, sorry, but
I have included the test cases for the "serial construct".
Best regards,
Frederik
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-libgomp-Fortran-Fix-OpenACC-gang-reduction-on-an-orp.patch --]
[-- Type: text/x-diff, Size: 9517 bytes --]
From 7c10ae450b95495dda362cb66770bb78b546592e Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
Date: Mon, 20 Jul 2020 11:24:21 +0200
Subject: [PATCH] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan
loop" error message
According to the OpenACC standard version 2.5 and later, reductions on
orphaned gang loops are explicitly disallowed (cf. section "Changes
from Version 2.0 to 2.5"). A loop is "orphaned" if it is not
lexically contained in a compute construct (cf. section "Loop
construct" of the OpenACC standard), i.e. in either a "parallel", a
"serial", or a "kernels" construct.
This commit fixes the check for reductions on orphaned gang loops in
the Fortran frontend which (in contrast to the C, C++ frontends)
erroneously rejects reductions on gang loops that are contained in
"kernels" constructs.
2020-07-20 Frederik Harwath <frederik@codesourcery.com>
gcc/fortran/
* openmp.c (oacc_is_parallel_or_serial): Removed function.
(oacc_is_kernels): New function.
(oacc_is_compute_construct): New 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-2.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-2.c: Likewise for C and C++.
---
gcc/fortran/ChangeLog | 9 ++
gcc/fortran/openmp.c | 13 ++-
gcc/testsuite/ChangeLog | 7 ++
.../c-c++-common/goacc/orphan-reductions-2.c | 103 ++++++++++++++++++
.../gfortran.dg/goacc/orphan-reductions-2.f90 | 87 +++++++++++++++
5 files changed, 216 insertions(+), 3 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index e86279cb647..5a1f81c286e 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,12 @@
+2020-07-20 Frederik Harwath <frederik@codesourcery.com>
+
+ * openmp.c (oacc_is_parallel_or_serial): Removed function.
+ (oacc_is_kernels): New function.
+ (oacc_is_compute_construct): New 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.
+
2020-07-08 Harald Anlauf <anlauf@gmx.de>
Backported from master:
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index ab68e9f2173..706933c869a 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5927,9 +5927,16 @@ oacc_is_serial (gfc_code *code)
}
static bool
-oacc_is_parallel_or_serial (gfc_code *code)
+oacc_is_kernels (gfc_code *code)
{
- return oacc_is_parallel (code) || oacc_is_serial (code);
+ return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP;
+}
+
+static bool
+oacc_is_compute_construct (gfc_code *code)
+{
+ return oacc_is_parallel (code) || oacc_is_serial (code)
+ || oacc_is_kernels (code);
}
static gfc_statement
@@ -6223,7 +6230,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))
+ 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/ChangeLog b/gcc/testsuite/ChangeLog
index 59e6c93b07a..fa1937a4ea2 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,10 @@
+2020-07-20 Frederik Harwath <frederik@codesourcery.com>
+
+ * gfortran.dg/goacc/orphan-reductions-2.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-2.c: Likewise for C and C++.
+
2020-07-12 Jakub Jelinek <jakub@redhat.com>
Backported from master:
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..d30321710dd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c
@@ -0,0 +1,103 @@
+/* Verify that the error message for gang reduction on orphaned OpenACC loops
+ is not reported for non-orphaned loops. */
+
+#include <assert.h>
+
+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-2.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
new file mode 100644
index 00000000000..6ad38039696
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
@@ -0,0 +1,87 @@
+! Verify that the error message for gang reductions on orphaned OpenACC loops
+! is not reported for non-orphaned loops.
+
+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.17.1
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] Make OpenACC orphan gang reductions errors
2017-05-02 1:32 [gomp4] Make OpenACC orphan gang reductions errors Cesar Philippidis
[not found] ` <xu8f8sfv4shz.fsf@harwath.name>
@ 2021-11-30 12:05 ` Thomas Schwinge
2021-11-30 12:10 ` Thomas Schwinge
2021-11-30 12:13 ` Thomas Schwinge
3 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2021-11-30 12:05 UTC (permalink / raw)
To: gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 1208 bytes --]
Hi!
On 2017-05-01T18:27:59-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> 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.
... which I've now adapted (with several things to be fixed in follow-up
commits) and pushed to master branch in
commit 2b7dac2c0dcb087da9e4018943c023c0678234a3
"Make OpenACC orphan gang reductions errors", see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Make-OpenACC-orphan-gang-reductions-errors.patch --]
[-- Type: text/x-diff, Size: 35061 bytes --]
From 2b7dac2c0dcb087da9e4018943c023c0678234a3 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Mon, 1 May 2017 18:27:59 -0700
Subject: [PATCH] Make OpenACC orphan gang reductions errors
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.
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 (oacc_is_parallel, oacc_is_kernels): New 'static'
functions.
(resolve_oacc_loop_blocks): Emit an error on orphan OpenACC gang
reductions.
gcc/
* omp-general.h (enum oacc_loop_flags): Add OLF_REDUCTION enum.
* omp-low.c (lower_oacc_head_mark): Use it to mark OpenACC
reductions.
* omp-offload.c (oacc_loop_auto_partitions): Don't assign gang
level parallelism to orphan reductions.
gcc/testsuite/
* c-c++-common/goacc/nested-reductions-1-routine.c: Adjust.
* c-c++-common/goacc/nested-reductions-2-routine.c: Likewise.
* gcc.dg/goacc/loop-processing-1.c: Likewise.
* gfortran.dg/goacc/nested-reductions-1-routine.f90: Likewise.
* gfortran.dg/goacc/nested-reductions-2-routine.f90: Likewise.
* c-c++-common/goacc/orphan-reductions-1.c: New test.
* c-c++-common/goacc/orphan-reductions-2.c: New test.
* gfortran.dg/goacc/orphan-reductions-1.f90: New test.
* gfortran.dg/goacc/orphan-reductions-2.f90: New test.
libgomp/
* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Temporarily
skip.
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
---
gcc/c/c-typeck.c | 8 +
gcc/cp/semantics.c | 8 +
gcc/fortran/openmp.c | 24 ++
gcc/omp-general.h | 3 +-
gcc/omp-low.c | 4 +
gcc/omp-offload.c | 7 +
.../goacc/nested-reductions-1-routine.c | 3 +
.../goacc/nested-reductions-2-routine.c | 9 +
.../c-c++-common/goacc/orphan-reductions-1.c | 56 +++++
.../c-c++-common/goacc/orphan-reductions-2.c | 87 ++++++++
.../gcc.dg/goacc/loop-processing-1.c | 2 +-
.../goacc/nested-reductions-1-routine.f90 | 3 +
.../goacc/nested-reductions-2-routine.f90 | 9 +
.../gfortran.dg/goacc/orphan-reductions-1.f90 | 206 ++++++++++++++++++
.../gfortran.dg/goacc/orphan-reductions-2.f90 | 89 ++++++++
.../libgomp.oacc-fortran/parallel-dims.f90 | 1 +
16 files changed, 517 insertions(+), 2 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/c/c-typeck.c b/gcc/c/c-typeck.c
index 7524304f2bd..a025740e618 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14135,6 +14135,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 && oacc_get_fn_attrib (current_function_decl)
+ && omp_find_clause (clauses, OMP_CLAUSE_GANG))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "gang reduction on an orphan loop");
+ remove = true;
+ break;
+ }
if (reduction_seen == 0)
reduction_seen = OMP_CLAUSE_REDUCTION_INSCAN (c) ? -1 : 1;
else if (reduction_seen != -2
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index cd1956497f8..c84caf43251 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6667,6 +6667,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 && oacc_get_fn_attrib (current_function_decl)
+ && omp_find_clause (clauses, OMP_CLAUSE_GANG))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "gang reduction on an orphan loop");
+ remove = true;
+ break;
+ }
if (reduction_seen == 0)
reduction_seen = OMP_CLAUSE_REDUCTION_INSCAN (c) ? -1 : 1;
else if (reduction_seen != -2
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index d120be81467..4fa38691c01 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -8322,6 +8322,17 @@ resolve_omp_do (gfc_code *code)
}
}
+static bool
+oacc_is_parallel (gfc_code *code)
+{
+ return code->op == EXEC_OACC_PARALLEL || code->op == EXEC_OACC_PARALLEL_LOOP;
+}
+
+static bool
+oacc_is_kernels (gfc_code *code)
+{
+ return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP;
+}
static gfc_statement
omp_code_to_statement (gfc_code *code)
@@ -8625,6 +8636,19 @@ resolve_oacc_loop_blocks (gfc_code *code)
if (!oacc_is_loop (code))
return;
+ if (code->op == EXEC_OACC_LOOP
+ && code->ext.omp_clauses->lists[OMP_LIST_REDUCTION]
+ && code->ext.omp_clauses->gang)
+ {
+ fortran_omp_context *c;
+ 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->tile_list && code->ext.omp_clauses->gang
&& code->ext.omp_clauses->worker && code->ext.omp_clauses->vector)
gfc_error ("Tiled loop cannot be parallelized across gangs, workers and "
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 8fe744c6a7a..a0c7c71148c 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -32,9 +32,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),
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 63a47f62d08..de3a26e08fc 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -8271,6 +8271,10 @@ lower_oacc_head_mark (location_t loc, tree ddvar, tree clauses,
tag |= OLF_TILE;
break;
+ case OMP_CLAUSE_REDUCTION:
+ tag |= OLF_REDUCTION;
+ break;
+
default:
continue;
}
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 833f7ddea58..0aec26b04e7 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1611,6 +1611,13 @@ oacc_loop_auto_partitions (oacc_loop *loop, unsigned outer_mask,
non-innermost available level. */
unsigned this_mask = GOMP_DIM_MASK (GOMP_DIM_GANG);
+ /* Orphan reductions cannot have gang partitioning. */
+ if ((loop->flags & OLF_REDUCTION)
+ && oacc_get_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/c-c++-common/goacc/nested-reductions-1-routine.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-routine.c
index 83d39950295..9e34614eb15 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-routine.c
@@ -44,6 +44,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(+:sum)
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
@@ -53,12 +54,14 @@ void acc_routine (void)
for (i = 0; i < 10; i++)
{
#pragma acc loop reduction(+:sum)
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(-:diff)
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(-:diff)
for (k = 0; k < 10; k++)
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c
index 5988d509bec..9bd79dea4cf 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-routine.c
@@ -11,6 +11,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
@@ -19,6 +20,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
for (k = 0; k < 10; k++)
#pragma acc loop reduction(+:sum)
@@ -28,6 +30,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
// { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -39,6 +42,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
for (k = 0; k < 10; k++)
@@ -47,6 +51,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(-:sum)
for (k = 0; k < 10; k++)
@@ -55,6 +60,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
// { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -66,6 +72,7 @@ void acc_routine (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." })
// { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -78,12 +85,14 @@ void acc_routine (void)
for (i = 0; i < 10; i++)
{
#pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(-:diff)
for (k = 0; k < 10; k++)
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..d2fec108214
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c
@@ -0,0 +1,56 @@
+/* Test orphan reductions. */
+
+/* { dg-do compile } */
+
+#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..941e5c6126a
--- /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-optimized-omp" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } */
+
+#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-optimized "assigned OpenACC worker vector loop 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-optimized "assigned OpenACC worker loop parallelism" } */
+ for (i = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC vector loop 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-optimized "assigned OpenACC worker loop parallelism" } */
+ for (i = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC seq loop parallelism" } */
+ /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */
+ for (j = 0; j < 100; j++)
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC vector loop 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-optimized "assigned OpenACC gang vector loop parallelism" } */
+ for (i = 0; i < 100; i++)
+ sum++;
+ }
+
+#pragma acc parallel copy (sum)
+ {
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC gang worker loop parallelism" } */
+ for (i = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC vector loop parallelism" } */
+ for (j = 0; j < 100; j++)
+ sum++;
+ }
+
+#pragma acc parallel copy (sum)
+ {
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC gang loop parallelism" } */
+ for (i = 0; i < 100; i++)
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC worker loop parallelism" } */
+ for (j = 0; j < 100; j++)
+#pragma acc loop reduction (+:sum) /* { dg-optimized "assigned OpenACC vector loop parallelism" } */
+ for (k = 0; k < 100; k++)
+ sum++;
+ }
+
+ return sum;
+}
diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
index 78b9aed89be..f6e25151e1e 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\);} "oaccloops" } } */
+/* { 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\);} "oaccloops" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90
index 17a586152c7..e8264114714 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-routine.f90
@@ -59,6 +59,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -70,6 +71,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum) reduction(-:diff)
do i = 1, 10
!$acc loop reduction(+:sum)
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -78,6 +80,7 @@ subroutine acc_routine ()
end do
!$acc loop reduction(-:diff)
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(-:diff)
do k = 1, 10
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90
index cc7802ecd10..98b1aa641c0 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-routine.f90
@@ -10,6 +10,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -21,6 +22,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
do k = 1, 10
!$acc loop reduction(+:sum)
@@ -34,6 +36,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -49,6 +52,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do k = 1, 10
@@ -60,6 +64,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(-:sum)
do k = 1, 10
@@ -71,6 +76,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -86,6 +92,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
@@ -101,6 +108,7 @@ subroutine acc_routine ()
!$acc loop reduction(+:sum) reduction(-:diff)
do i = 1, 10
!$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
@@ -109,6 +117,7 @@ subroutine acc_routine ()
end do
!$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
do j = 1, 10
!$acc loop reduction(-:diff)
do k = 1, 10
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..7ff0a57e620
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90
@@ -0,0 +1,89 @@
+! Ensure that the middle end does not assign gang level parallelism to
+! orphan loop containing reductions.
+
+! { dg-do compile }
+! { dg-additional-options "-fopt-info-optimized-omp" }
+! { dg-additional-options "-Wopenacc-parallelism" }
+
+subroutine s1 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" }
+ implicit none
+ !$acc routine gang
+ integer i, sum
+
+ sum = 0
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker vector loop 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
+
+ sum = 0
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker loop parallelism" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC vector loop 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
+
+ sum = 0
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker loop parallelism" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC seq loop parallelism" }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC vector loop 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
+
+ sum = 0
+ !$acc parallel copy(sum)
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC gang vector loop parallelism" }
+ do i = 1, 10
+ sum = sum + 1
+ end do
+ !$acc end parallel
+
+ !$acc parallel copy(sum)
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC gang worker loop parallelism" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC vector loop parallelism" }
+ do j = 1, 10
+ sum = sum + 1
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel copy(sum)
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC gang loop parallelism" }
+ do i = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC worker loop parallelism" }
+ do j = 1, 10
+ !$acc loop reduction (+:sum) ! { dg-optimized "assigned OpenACC vector loop parallelism" }
+ do k = 1, 10
+ sum = sum + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+end subroutine s4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
index fad3d9d6a80..80d64030414 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
@@ -3,6 +3,7 @@
! { dg-additional-sources parallel-dims-aux.c }
! { dg-do run }
+ ! { dg-skip-if TODO { *-*-* } }
! { dg-prune-output "command-line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" }
! { dg-additional-options "-fopt-info-note-omp" }
--
2.33.0
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] Make OpenACC orphan gang reductions errors
2017-05-02 1:32 [gomp4] Make OpenACC orphan gang reductions errors Cesar Philippidis
[not found] ` <xu8f8sfv4shz.fsf@harwath.name>
2021-11-30 12:05 ` [gomp4] Make OpenACC orphan gang reductions errors Thomas Schwinge
@ 2021-11-30 12:10 ` Thomas Schwinge
2021-11-30 12:13 ` Thomas Schwinge
3 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2021-11-30 12:10 UTC (permalink / raw)
To: gcc-patches, fortran; +Cc: Kwok Cheung Yeung
[-- Attachment #1: Type: text/plain, Size: 1260 bytes --]
Hi!
On 2017-05-01T18:27:59-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -6090,6 +6090,18 @@ resolve_oacc_loop_blocks (gfc_code *code)
> + 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);
> + }
To avoid erroneous diagnostics, we also need to handle the OpenACC
'serial' construct here. I've adapted Kwok's relevant patch, and pushed
to master branch commit f1a58ab0db20c0862e8b5039bd448fc8c9799cac
"[OpenACC] Allow gang reductions inside serial constructs", see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-Allow-gang-reductions-inside-serial-construc.patch --]
[-- Type: text/x-diff, Size: 2462 bytes --]
From f1a58ab0db20c0862e8b5039bd448fc8c9799cac Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcy@codesourcery.com>
Date: Fri, 13 Mar 2020 11:13:49 -0700
Subject: [PATCH] [OpenACC] Allow gang reductions inside serial constructs
... fixing a regression introduced in the preceding
commit 2b7dac2c0dcb087da9e4018943c023c0678234a3
"Make OpenACC orphan gang reductions errors".
gcc/fortran/
* openmp.c (oacc_is_serial, oacc_is_parallel_or_serial): New.
(resolve_oacc_loop_blocks): Use oacc_is_parallel_or_serial instead of
oacc_is_parallel.
libgomp/
* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Remove
temporary skip.
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
---
gcc/fortran/openmp.c | 14 +++++++++++++-
.../libgomp.oacc-fortran/parallel-dims.f90 | 1 -
2 files changed, 13 insertions(+), 2 deletions(-)
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 4fa38691c01..b4100577e51 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -8334,6 +8334,18 @@ oacc_is_kernels (gfc_code *code)
return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP;
}
+static bool
+oacc_is_serial (gfc_code *code)
+{
+ return code->op == EXEC_OACC_SERIAL || code->op == EXEC_OACC_SERIAL_LOOP;
+}
+
+static bool
+oacc_is_parallel_or_serial (gfc_code *code)
+{
+ return oacc_is_parallel (code) || oacc_is_serial (code);
+}
+
static gfc_statement
omp_code_to_statement (gfc_code *code)
{
@@ -8644,7 +8656,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 (c->code)
+ if (c == NULL || !(oacc_is_parallel_or_serial (c->code)
|| oacc_is_kernels (c->code)))
gfc_error ("gang reduction on an orphan loop at %L", &code->loc);
}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
index 80d64030414..fad3d9d6a80 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
@@ -3,7 +3,6 @@
! { dg-additional-sources parallel-dims-aux.c }
! { dg-do run }
- ! { dg-skip-if TODO { *-*-* } }
! { dg-prune-output "command-line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" }
! { dg-additional-options "-fopt-info-note-omp" }
--
2.33.0
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message
2020-07-20 10:26 ` Frederik Harwath
@ 2021-11-30 12:12 ` Thomas Schwinge
0 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2021-11-30 12:12 UTC (permalink / raw)
To: fortran, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 922 bytes --]
Hi!
On 2020-07-20T12:26:48+0200, Frederik Harwath <frederik@codesourcery.com> wrote:
> Thomas Schwinge <thomas@codesourcery.com> writes:
>>> Can I include the patch in OG10?
> This has been delayed a bit by my vacation, but I have now committed
> the patch.
>> (Ideally, we'd also test 'serial' construct in addition to 'kernels',
>> 'parallel'
> I have included the test cases for the "serial construct".
I've adapted the remaining relevant changes and pushed to master branch
commit c4f4c60457d1657cbd72015de3d818eb6462a0e9
'Re OpenACC "gang reduction on an orphan loop" error message', see
attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Re-OpenACC-gang-reduction-on-an-orphan-loop-error-me.patch --]
[-- Type: text/x-diff, Size: 7726 bytes --]
From c4f4c60457d1657cbd72015de3d818eb6462a0e9 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
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 <thomas@codesourcery.com>
---
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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] Make OpenACC orphan gang reductions errors
2017-05-02 1:32 [gomp4] Make OpenACC orphan gang reductions errors Cesar Philippidis
` (2 preceding siblings ...)
2021-11-30 12:10 ` Thomas Schwinge
@ 2021-11-30 12:13 ` Thomas Schwinge
3 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2021-11-30 12:13 UTC (permalink / raw)
To: gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 937 bytes --]
Hi!
On 2017-05-01T18:27:59-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> 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.
As a follow-up, I've pushed to master branch
commit 77d24d43644909852998043335b5a0e09d1e8f02
'Consolidate OpenACC "gang reduction on an orphan loop" checking',
see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Consolidate-OpenACC-gang-reduction-on-an-orphan-loop.patch --]
[-- Type: text/x-diff, Size: 7003 bytes --]
From 77d24d43644909852998043335b5a0e09d1e8f02 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 26 Nov 2021 12:29:26 +0100
Subject: [PATCH] Consolidate OpenACC "gang reduction on an orphan loop"
checking
No need to implement separately in all front ends what we may implement in the
middle end, once for all.
Follow-up to preceding commit 2b7dac2c0dcb087da9e4018943c023c0678234a3
"Make OpenACC orphan gang reductions errors".
gcc/
* omp-offload.c (oacc_loop_process): Implement "gang reduction on
an orphan loop" checking.
gcc/c/
* c-typeck.c (c_finish_omp_clauses): Remove "gang reduction on an
orphan loop" checking.
gcc/cp/
* semantics.c (finish_omp_clauses): Remove "gang reduction on an
orphan loop" checking.
gcc/fortran/
* openmp.c (resolve_oacc_loop_blocks): Remove "gang reduction on
an orphan loop" checking.
(oacc_is_parallel, oacc_is_kernels, oacc_is_serial)
(oacc_is_compute_construct): Remove.
gcc/testsuite/
* gfortran.dg/goacc/orphan-reductions-1.f90: Adjust.
---
gcc/c/c-typeck.c | 8 ----
gcc/cp/semantics.c | 8 ----
gcc/fortran/openmp.c | 37 -------------------
gcc/omp-offload.c | 20 ++++++++--
.../gfortran.dg/goacc/orphan-reductions-1.f90 | 8 ++--
5 files changed, 20 insertions(+), 61 deletions(-)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index a025740e618..7524304f2bd 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14135,14 +14135,6 @@ 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 && oacc_get_fn_attrib (current_function_decl)
- && omp_find_clause (clauses, OMP_CLAUSE_GANG))
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "gang reduction on an orphan loop");
- remove = true;
- break;
- }
if (reduction_seen == 0)
reduction_seen = OMP_CLAUSE_REDUCTION_INSCAN (c) ? -1 : 1;
else if (reduction_seen != -2
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index c84caf43251..cd1956497f8 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6667,14 +6667,6 @@ 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 && oacc_get_fn_attrib (current_function_decl)
- && omp_find_clause (clauses, OMP_CLAUSE_GANG))
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "gang reduction on an orphan loop");
- remove = true;
- break;
- }
if (reduction_seen == 0)
reduction_seen = OMP_CLAUSE_REDUCTION_INSCAN (c) ? -1 : 1;
else if (reduction_seen != -2
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 7950c7fb43d..d120be81467 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -8322,31 +8322,6 @@ resolve_omp_do (gfc_code *code)
}
}
-static bool
-oacc_is_parallel (gfc_code *code)
-{
- return code->op == EXEC_OACC_PARALLEL || code->op == EXEC_OACC_PARALLEL_LOOP;
-}
-
-static bool
-oacc_is_kernels (gfc_code *code)
-{
- return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP;
-}
-
-static bool
-oacc_is_serial (gfc_code *code)
-{
- return code->op == EXEC_OACC_SERIAL || code->op == EXEC_OACC_SERIAL_LOOP;
-}
-
-static bool
-oacc_is_compute_construct (gfc_code *code)
-{
- return (oacc_is_parallel (code)
- || oacc_is_kernels (code)
- || oacc_is_serial (code));
-}
static gfc_statement
omp_code_to_statement (gfc_code *code)
@@ -8650,18 +8625,6 @@ resolve_oacc_loop_blocks (gfc_code *code)
if (!oacc_is_loop (code))
return;
- if (code->op == EXEC_OACC_LOOP
- && code->ext.omp_clauses->lists[OMP_LIST_REDUCTION]
- && code->ext.omp_clauses->gang)
- {
- fortran_omp_context *c;
- for (c = omp_current_ctx; c; c = c->previous)
- if (!oacc_is_loop (c->code))
- break;
- if (c == NULL || !(oacc_is_compute_construct (c->code)))
- gfc_error ("gang reduction on an orphan loop at %L", &code->loc);
- }
-
if (code->ext.omp_clauses->tile_list && code->ext.omp_clauses->gang
&& code->ext.omp_clauses->worker && code->ext.omp_clauses->vector)
gfc_error ("Tiled loop cannot be parallelized across gangs, workers and "
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 0aec26b04e7..5110a424584 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1380,10 +1380,10 @@ oacc_loop_xform_head_tail (gcall *from, int level)
partitioning level etc. */
static void
-oacc_loop_process (oacc_loop *loop)
+oacc_loop_process (oacc_loop *loop, int fn_level)
{
if (loop->child)
- oacc_loop_process (loop->child);
+ oacc_loop_process (loop->child, fn_level);
if (loop->mask && !loop->routine)
{
@@ -1432,7 +1432,19 @@ oacc_loop_process (oacc_loop *loop)
}
if (loop->sibling)
- oacc_loop_process (loop->sibling);
+ oacc_loop_process (loop->sibling, fn_level);
+
+
+ /* OpenACC 2.6, 2.9.11. "reduction clause" places a restriction such that
+ "The 'reduction' clause may not be specified on an orphaned 'loop'
+ construct with the 'gang' clause, or on an orphaned 'loop' construct that
+ will generate gang parallelism in a procedure that is compiled with the
+ 'routine gang' clause." */
+ if (fn_level == GOMP_DIM_GANG
+ && (loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG))
+ && (loop->flags & OLF_REDUCTION))
+ error_at (loop->loc,
+ "gang reduction on an orphan loop");
}
/* Walk the OpenACC loop heirarchy checking and assigning the
@@ -2072,7 +2084,7 @@ execute_oacc_loop_designation ()
if (is_oacc_parallel_kernels_gang_single)
gcc_checking_assert (dims[GOMP_DIM_GANG] == 1);
- oacc_loop_process (loops);
+ oacc_loop_process (loops, fn_level);
if (dump_file)
{
fprintf (dump_file, "OpenACC loops\n");
diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90
index c7fcc9d4ac5..464dee1260a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90
@@ -30,7 +30,7 @@ end subroutine s1
subroutine s2
implicit none
- !$acc routine worker
+ !$acc routine gang
integer, parameter :: n = 100
integer :: i, j, sum
@@ -80,7 +80,7 @@ end function f1
integer function f2 ()
implicit none
- !$acc routine worker
+ !$acc routine gang
integer, parameter :: n = 100
integer :: i, j, sum
@@ -132,7 +132,7 @@ contains
subroutine s4
implicit none
- !$acc routine worker
+ !$acc routine gang
integer, parameter :: n = 100
integer :: i, j, sum
@@ -182,7 +182,7 @@ contains
integer function f4 ()
implicit none
- !$acc routine worker
+ !$acc routine gang
integer, parameter :: n = 100
integer :: i, j, sum
--
2.33.0
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2021-11-30 12:14 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2017-05-02 1:32 [gomp4] Make OpenACC orphan gang reductions errors Cesar Philippidis
[not found] ` <xu8f8sfv4shz.fsf@harwath.name>
2020-07-07 10:42 ` [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message Thomas Schwinge
2020-07-07 12:39 ` Frederik Harwath
2020-07-20 10:26 ` Frederik Harwath
2021-11-30 12:12 ` Thomas Schwinge
2021-11-30 12:05 ` [gomp4] Make OpenACC orphan gang reductions errors Thomas Schwinge
2021-11-30 12:10 ` Thomas Schwinge
2021-11-30 12:13 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).