Hi Cesar! On Wed, 22 Mar 2017 06:40:03 -0700, Cesar Philippidis <ce...@codesourcery.com> wrote: > In addition to resolving the memory leak involving OpenACC delare > clauses, this patch also corrects an ICE involving VLA variables as data > clause arguments used in acc declare constructs. More details can be > found here <https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>. > > Is this OK for trunk?
(Got committed in r246381.) > PR c++/80029 > > gcc/ > * gimplify.c (is_oacc_declared): New function. > (oacc_default_clause): Use it to set default flags for acc declared > variables inside parallel regions. > (gimplify_scan_omp_clauses): Strip firstprivate pointers for acc > declared variables. > (gimplify_oacc_declare): Gimplify the declare clauses. Add the > declare attribute to any decl as necessary. > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -6787,6 +6787,16 @@ device_resident_p (tree decl) > return false; > } > > +/* Return true if DECL has an ACC DECLARE attribute. */ > + > +static bool > +is_oacc_declared (tree decl) > +{ Adding here a "gcc_assert (TREE_CODE (decl) != MEM_REF);", that one never triggers (at least given the current test coverage), so... > + tree t = TREE_CODE (decl) == MEM_REF ? TREE_OPERAND (decl, 0) : decl; ... it would seem that this code can be simplified -- or under which circumstances should we get a MEM_REF here? The call from gimplify_oacc_declare itself does the tree operand 0 indirection, and the call from oacc_default_clause shouldn't be a MEM_REF, I would think. > + tree declared = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES > (t)); > + return declared != NULL_TREE; > +} > @@ -6887,6 +6897,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree > decl, unsigned flags) > { > const char *rkind; > bool on_device = false; > + bool declared = is_oacc_declared (decl); > tree type = TREE_TYPE (decl); > > if (lang_hooks.decls.omp_privatize_by_reference (decl)) > @@ -6917,7 +6928,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree > decl, unsigned flags) > > case ORT_ACC_PARALLEL: > { > - if (on_device || AGGREGATE_TYPE_P (type)) > + if (on_device || AGGREGATE_TYPE_P (type) || declared) > /* Aggregates default to 'present_or_copy'. */ > flags |= GOVD_MAP; > else Isn't this new "declared" logic doing a thing very similar to the existing "on_device" logic? Should that be combined/cleaned up? Why doesn't the same "declared"/"on_device" logic also apply to ORT_ACC_KERNELS? Grüße Thomas > @@ -7346,6 +7357,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > case OMP_TARGET_DATA: > case OMP_TARGET_ENTER_DATA: > case OMP_TARGET_EXIT_DATA: > + case OACC_DECLARE: > case OACC_HOST_DATA: > ctx->target_firstprivatize_array_bases = true; > default: > @@ -9231,18 +9243,26 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq > *pre_p) > { > tree expr = *expr_p; > gomp_target *stmt; > - tree clauses, t; > + tree clauses, t, decl; > > clauses = OACC_DECLARE_CLAUSES (expr); > > gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE); > + gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE); > > for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) > { > - tree decl = OMP_CLAUSE_DECL (t); > + decl = OMP_CLAUSE_DECL (t); > > if (TREE_CODE (decl) == MEM_REF) > - continue; > + decl = TREE_OPERAND (decl, 0); > + > + if (VAR_P (decl) && !is_oacc_declared (decl)) > + { > + tree attr = get_identifier ("oacc declare target"); > + DECL_ATTRIBUTES (decl) = tree_cons (attr, NULL_TREE, > + DECL_ATTRIBUTES (decl)); > + } > > if (VAR_P (decl) > && !is_global_var (decl) > @@ -9258,7 +9278,8 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) > } > } > > - omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); > + if (gimplify_omp_ctxp) > + omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); > } > > stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c > b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c > new file mode 100644 > index 0000000..3ea148e > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c > @@ -0,0 +1,25 @@ > +/* Verify that acc declare accept VLA variables. */ > + > +#include <assert.h> > + > +int > +main () > +{ > + int N = 1000; > + int i, A[N]; > +#pragma acc declare copy(A) > + > + for (i = 0; i < N; i++) > + A[i] = -i; > + > +#pragma acc kernels > + for (i = 0; i < N; i++) > + A[i] = i; > + > +#pragma acc update host(A) > + > + for (i = 0; i < N; i++) > + assert (A[i] == i); > + > + return 0; > +}