This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [PATCH] Fix PR80029
- From: Thomas Schwinge <thomas at codesourcery dot com>
- To: Cesar Philippidis <cesar at codesourcery dot com>
- Cc: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, Jakub Jelinek <jakub at redhat dot com>
- Date: Fri, 7 Apr 2017 11:52:47 +0200
- Subject: Re: [PATCH] Fix PR80029
- Authentication-results: sourceware.org; auth=none
- References: <bf3bc2cb-f7b2-b10f-ed06-f257724cbb52@codesourcery.com> <91b16e24-bd33-63db-dbbf-37151c0d584a@codesourcery.com> <278ce4b7-cc02-453b-10fc-8094e9078d9a@codesourcery.com>
Hi Cesar!
On Wed, 22 Mar 2017 06:40:03 -0700, Cesar Philippidis <cesar@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;
> +}