This patch teaches the c and c++ front ends not to add a 'copy' clause for each non-loop OpenACC reduction. Without this patch, this construct would error with a duplicate mapping for 'sum':
#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) While the intent behind adding a 'copy' clause for the reduction variable was good, it is inconsistent with the behavior of the other data clauses in the OpenACC spec. Without an explicit copy, 'sum' should probably be transferred to the accelerator as firstprivate. Anyway, it's probably better to correct this behavior in the gimplifier later on if necessary. I'll apply this patch to gomp-4_0-branch shortly. Cesar
2015-08-14 Cesar Philippidis <ce...@codesourcery.com> gcc/ * c/c-typeck.c (c_finish_omp_clauses): Permit variables to appear in both OpenACC data and reduction clauses. * cp/semantics.c (finish_omp_clauses): Likewise. gcc/testsuite/ * c-c++-common/goacc/parallel-reduction.c: New test. diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index ed748c8..fe54799f 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12059,11 +12059,10 @@ tree c_finish_omp_clauses (tree clauses, bool oacc) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, oacc_data_head; + bitmap_head aligned_head, oacc_data_head, oacc_reduction_head; tree c, t, *pc; bool branch_seen = false; bool copyprivate_seen = false; - bool oacc_data = false; tree *nowait_clause = NULL; bitmap_obstack_initialize (NULL); @@ -12072,12 +12071,15 @@ c_finish_omp_clauses (tree clauses, bool oacc) bitmap_initialize (&lastprivate_head, &bitmap_default_obstack); bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&oacc_data_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { bool remove = false; bool need_complete = false; bool need_implicitly_determined = false; + bool oacc_data = false; + bool reduction = false; switch (OMP_CLAUSE_CODE (c)) { @@ -12095,8 +12097,8 @@ c_finish_omp_clauses (tree clauses, bool oacc) goto check_dup_generic; case OMP_CLAUSE_REDUCTION: - need_implicitly_determined = true; - oacc_data = false; + need_implicitly_determined = !oacc; + reduction = true; t = OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c) == NULL_TREE && (FLOAT_TYPE_P (TREE_TYPE (t)) @@ -12312,12 +12314,23 @@ c_finish_omp_clauses (tree clauses, bool oacc) else bitmap_set_bit (&oacc_data_head, DECL_UID (t)); } + else if (reduction) + { + if (oacc && bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears in multiple reduction clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_reduction_head, DECL_UID (t)); + } else { if (bitmap_bit_p (&generic_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), - "%qE appears more than once in data clauses", t); + "%qE appears more than one non-data clause", t); remove = true; } else diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 0b13ca2..cf1790c 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5281,11 +5281,10 @@ tree finish_omp_clauses (tree clauses, bool oacc) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, oacc_data_head; + bitmap_head aligned_head, oacc_data_head, oacc_reduction_head; tree c, t, *pc; bool branch_seen = false; bool copyprivate_seen = false; - bool oacc_data = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -5293,10 +5292,13 @@ finish_omp_clauses (tree clauses, bool oacc) bitmap_initialize (&lastprivate_head, &bitmap_default_obstack); bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&oacc_data_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { bool remove = false; + bool oacc_data = false; + bool reduction = false; switch (OMP_CLAUSE_CODE (c)) { @@ -5314,6 +5316,7 @@ finish_omp_clauses (tree clauses, bool oacc) if (oacc) { oacc_data = false; + reduction = true; goto check_dup_oacc; } else @@ -5426,6 +5429,17 @@ finish_omp_clauses (tree clauses, bool oacc) else bitmap_set_bit (&oacc_data_head, DECL_UID (t)); } + else if (reduction) + { + if (oacc && bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears in multiple reduction clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_reduction_head, DECL_UID (t)); + } else { if (bitmap_bit_p (&generic_head, DECL_UID (t))) diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c new file mode 100644 index 0000000..debed55 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c @@ -0,0 +1,13 @@ +int +main () +{ + int sum = 0; + +#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) + { + int v = 5; + sum += 10 + v; + } + + return sum; +}