2016-04-04 Cesar Philippidis <cesar@codesourcery.com>
gcc/
* gimplify.c (gimplify_adjust_acc_parallel_reductions): New function.
(gimplify_omp_workshare): Call it. Add new data clauses for acc
parallel reductions as needed.
gcc/testsuite/
* c-c++-common/goacc/reduction-5.c: New test.
* c-c++-common/goacc/reduction-promotions.c: New test.
* gfortran.dg/goacc/reduction-3.f95: New test.
* gfortran.dg/goacc/reduction-promotions.f90: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70289.c: New test.
* testsuite/libgomp.oacc-c-c++-common/pr70373.c: New test.
* testsuite/libgomp.oacc-fortran/pr70289.f90: New test.
@@ -9484,6 +9484,108 @@ optimize_target_teams (tree target, gimple_seq *pre_p)
OMP_TARGET_CLAUSES (target) = c;
}
+/* OpenACC parallel reductions need a present_or_copy clause to ensure
+ that the original variable used in the reduction gets updated on
+ the host. This function scans CLAUSES for reductions and adds or
+ adjusts the data clauses as necessary. Any incompatible data clause
+ will be reported as a warning and promoted to present_or_copy. Any
+ private reduction will be treated as an error. This function
+ returns a list of new present_or_copy date clauses. */
+
+static tree
+gimplify_adjust_acc_parallel_reductions (tree *clauses)
+{
+ tree c, list = NULL_TREE;
+ hash_set<tree> *reduction_decls;
+ reduction_decls = new hash_set<tree>;
+
+ /* Scan 1: Construct a hash set with all of the reduction decls. */
+ for (c = *clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ reduction_decls->add (OMP_CLAUSE_DECL (c));
+ }
+
+ if (reduction_decls->elements () == 0)
+ goto cleanup;
+
+ /* Scan 2: Adjust the data clause for each reduction. */
+ for (c = *clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ int kind = -1;
+ tree decl;
+
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_MAP:
+ kind = OMP_CLAUSE_MAP_KIND (c);
+ case OMP_CLAUSE_PRIVATE:
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ decl = OMP_CLAUSE_DECL (c);
+
+ /* Reference variables always have a GOMP_MAP_ALLOC. Ignore it. */
+ if (POINTER_TYPE_P (TREE_TYPE (decl))
+ && kind == GOMP_MAP_ALLOC)
+ break;
+
+ if (!DECL_P (decl))
+ decl = TREE_OPERAND (decl, 0);
+ gcc_assert (DECL_P (decl));
+
+ if (!reduction_decls->contains (decl))
+ break;
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ {
+ if (!((kind & GOMP_MAP_TOFROM) == GOMP_MAP_TOFROM
+ || kind == GOMP_MAP_FORCE_PRESENT))
+ {
+ warning_at (OMP_CLAUSE_LOCATION (c), 0, "incompatible data "
+ "clause with reduction on %qE; promoting to "
+ "present_or_copy", DECL_NAME (decl));
+
+ OMP_CLAUSE_CODE (c) = OMP_CLAUSE_MAP;
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+ }
+ reduction_decls->remove (decl);
+ break;
+ }
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "invalid private reduction "
+ "on %qE", DECL_NAME (decl));
+ reduction_decls->remove (decl);
+ }
+ default:;
+ }
+ }
+
+ if (reduction_decls->elements () == 0)
+ goto cleanup;
+
+ /* Scan 3: Add a present_or_copy clause for any reduction variable which
+ doens't have a data clause already. */
+
+ for (hash_set<tree>::iterator iter = reduction_decls->begin ();
+ iter != reduction_decls->end (); ++iter)
+ {
+ tree decl = *iter;
+
+ tree nc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_TOFROM);
+ OMP_CLAUSE_DECL (nc) = decl;
+ TREE_CHAIN (nc) = list;
+ list = nc;
+ }
+
+ cleanup:
+ delete reduction_decls;
+
+ return list;
+}
+
/* Gimplify the gross structure of several OMP constructs. */
static void
@@ -9491,6 +9593,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
{
tree expr = *expr_p;
gimple *stmt;
+ tree acc_reductions = NULL_TREE;
gimple_seq body = NULL;
enum omp_region_type ort;
@@ -9508,6 +9611,8 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
break;
case OACC_PARALLEL:
ort = ORT_ACC_PARALLEL;
+ acc_reductions
+ = gimplify_adjust_acc_parallel_reductions (&OMP_CLAUSES (expr));
break;
case OACC_DATA:
ort = ORT_ACC_DATA;
@@ -9524,6 +9629,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
default:
gcc_unreachable ();
}
+
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
TREE_CODE (expr));
if (TREE_CODE (expr) == OMP_TARGET)
@@ -9606,6 +9712,41 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
gimplify_seq_add_stmt (pre_p, stmt);
*expr_p = NULL_TREE;
+
+ /* Finalize any parallel acc reductions. */
+ if (acc_reductions)
+ {
+ tree c, nc, t;
+ tree clauses = NULL_TREE;
+
+ c = nc = acc_reductions;
+
+ while (c)
+ {
+ nc = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = NULL_TREE;
+ lang_hooks.decls.omp_finish_clause (c, pre_p);
+
+ /* Find the last data clause introduced by omp_finish_decls. */
+ for (t = c; TREE_CHAIN (t); t = TREE_CHAIN (t))
+ ;
+
+ /* Update the chain of clauses. */
+ TREE_CHAIN (t) = clauses;
+ clauses = c;
+
+ c = nc;
+ }
+
+ /* Update the list of clauses in the gimple stmt. */
+ for (t = gimple_omp_target_clauses (stmt); OMP_CLAUSE_CHAIN (t);
+ t = OMP_CLAUSE_CHAIN (t))
+ ;
+
+ OMP_CLAUSE_CHAIN (t) = clauses;
+ }
+
+ return;
}
/* Gimplify the gross structure of OpenACC enter/exit data, update, and OpenMP
new file mode 100644
@@ -0,0 +1,16 @@
+/* Integer reductions. */
+
+#define n 1000
+
+int
+main(void)
+{
+ int v1;
+
+#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "appears more than once in data clauses" } */
+ ;
+#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "appears more than once in data clauses" } */
+ ;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,32 @@
+/* Integer reductions. */
+
+#define n 1000
+
+int
+main(void)
+{
+ int v1, v2;
+
+#pragma acc parallel reduction(+:v1,v2)
+ ;
+#pragma acc parallel reduction(+:v1,v2) copy(v1,v2)
+ ;
+#pragma acc parallel reduction(+:v1,v2) pcopy(v1,v2)
+ ;
+#pragma acc parallel reduction(+:v1,v2) present(v1,v2)
+ ;
+#pragma acc parallel reduction(+:v1,v2) copyin(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+ ;
+#pragma acc parallel reduction(+:v1,v2) pcopyin(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+ ;
+#pragma acc parallel reduction(+:v1,v2) copyout(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+ ;
+#pragma acc parallel reduction(+:v1,v2) pcopyout(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+ ;
+#pragma acc parallel reduction(+:v1,v2) create(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+ ;
+#pragma acc parallel reduction(+:v1,v2) pcreate(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+ ;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,10 @@
+! { dg-do compile }
+
+subroutine foo (ia1)
+integer :: i1, i2
+
+!$acc parallel reduction (+:i1) private(i1) ! { dg-error "invalid private reduction on .i1." }
+!$acc end parallel
+!$acc parallel reduction (+:i2) firstprivate(i2) ! { dg-error "invalid private reduction on .i2." }
+!$acc end parallel
+end subroutine foo
new file mode 100644
@@ -0,0 +1,45 @@
+! Ensure that each parallel reduction variable as a copy or pcopy
+! data clause.
+
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program test
+ implicit none
+ integer :: v1, v2
+
+ !$acc parallel reduction(+:v1,v2)
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) copy(v1,v2)
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) pcopy(v1,v2)
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) present(v1,v2)
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) copyin(v1,v2) ! { dg-warning "incompatible data clause" }
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) pcopyin(v1,v2) ! { dg-warning "incompatible data clause" }
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) copyout(v1,v2) ! { dg-warning "incompatible data clause" }
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) pcopyout(v1,v2) ! { dg-warning "incompatible data clause" }
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) create(v1,v2) ! { dg-warning "incompatible data clause" }
+ !$acc end parallel
+
+ !$acc parallel reduction(+:v1,v2) pcreate(v1,v2) ! { dg-warning "incompatible data clause" }
+ !$acc end parallel
+end program test
+! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } }
new file mode 100644
@@ -0,0 +1,13 @@
+int
+main ()
+{
+ int i;
+ static int temp;
+
+#pragma acc parallel reduction(+:temp)
+ {
+ temp++;
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,20 @@
+#define N 32
+
+int
+foo (unsigned int sum)
+{
+#pragma acc parallel reduction (+:sum)
+ {
+ sum;
+ }
+
+ return sum;
+}
+
+int
+main (void)
+{
+ unsigned int sum = 0;
+ foo (sum);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,20 @@
+program foo
+ implicit none
+ integer :: i
+ integer :: temp = 0
+ integer :: temp2 = 0
+
+ !$acc parallel
+ !$acc loop gang private(temp)
+ do i=1, 10000
+ temp = 0
+ enddo
+ !$acc end parallel
+
+ !$acc parallel reduction(+:temp2)
+ !$acc loop gang reduction(+:temp2)
+ do i=1, 10000
+ temp2 = 0
+ enddo
+ !$acc end parallel
+end program foo