On Thu, Apr 07, 2016 at 09:34:43PM -0700, Cesar Philippidis wrote:
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -5802,7 +5802,8 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree
> decl, unsigned int flags)
> flags |= GOVD_SEEN;
>
> n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
> - if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
> + if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0
> + && ctx->region_type != ORT_ACC_PARALLEL)
> {
> /* We shouldn't be re-adding the decl with the same data
> sharing class. */
Why?
> @@ -6557,6 +6558,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
> *pre_p,
> break;
> }
>
> +/* OpenACC parallel reductions need a present_or_copy clause to ensure
> + that the original variable used in the reduction gets updated on
> + the host. Scan the list of clauses for reduction so that any existing
> + data clause can be adjusted if necessary. */
> + if (region_type == ORT_ACC_PARALLEL)
> + {
> + for (c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
> + {
> + tree decl = NULL_TREE;
> +
> + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION)
> + continue;
> +
> + decl = OMP_CLAUSE_DECL (c);
> + omp_add_variable (ctx, decl, GOVD_REDUCTION);
> + }
> + }
> +
And this looks also wrong, why?
If I try under the debugger 3 cases:
void f1 (int sum)
{
#pragma acc parallel reduction(+:sum) present_or_copy(sum)
;
}
void f2 (int sum)
{
#pragma acc parallel present_or_copy(sum)
;
}
void f3 (int sum)
{
#pragma acc parallel reduction(+:sum)
;
}
then I see the loop that starts with the while below doing the right thing
already. In the first case you end up with
GOVD_SEEN | GOVD_EXPLICIT | GOVD_REDUCTION | GOVD_MAP
in the second with
GOVD_SEEN | GOVD_EXPLICIT | GOVD_MAP
and third one with
GOVD_SEEN | GOVD_EXPLICIT | GOVD_REDUCTION
That is where you IMHO should stop at the gimplify_scan_omp_clauses side,
so don't modify neither omp_add_variable nor gimplify_scan_omp_clauses
at all, and do everything else in gimplify_adjust_omp_clauses.
That function walks the explicit clauses and has all the info gathered
during gimplify_scan_omp_clauses available in the splay tree.
So, you can do all the checking there. Say on OMP_CLAUSE_REDUCTION
for the ORT_ACC_PARALLEL check the flags if they include GOVD_PRIVATE
or GOVD_FIRSTPRIVATE, if yes, complain. Also check if GOVD_MAP is included,
if not, add the extra OMP_CLAUSE_MAP tofrom.
And, on OMP_CLAUSE_MAP, check if GOVD_REDUCTION is set on ORT_ACC_PARALLEL,
and if yes, check if it is tofrom and complain otherwise.
> while ((c = *list_p) != NULL)
> {
> bool remove = false;
> @@ -6808,6 +6827,31 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
> *pre_p,
>
> case OMP_CLAUSE_MAP:
> decl = OMP_CLAUSE_DECL (c);
> + if (region_type == ORT_ACC_PARALLEL)
> + {
> + tree t = DECL_P (decl) ? decl : TREE_OPERAND (decl, 0);
> + splay_tree_node n = NULL;
> +
> + if (DECL_P (t))
> + n = splay_tree_lookup (ctx->variables, (splay_tree_key)t);
> +
> + if (n)
> + {
> + int kind = OMP_CLAUSE_MAP_KIND (c);
> +
> + OMP_CLAUSE_MAP_IN_REDUCTION(c) = 1;
> + if ((kind & GOMP_MAP_TOFROM) != GOMP_MAP_TOFROM
> + && kind != GOMP_MAP_FORCE_PRESENT
> + && kind != GOMP_MAP_POINTER)
> + {
> + warning_at (OMP_CLAUSE_LOCATION (c), 0,
> + "incompatible data clause with reduction "
> + "on %qE; promoting to present_or_copy",
> + DECL_NAME (t));
> + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
> + }
> + }
> + }
> if (error_operand_p (decl))
> remove = true;
> switch (code)
So the above is also wrong IMNSHO.
Jakub
