[PATCH][openacc] Fix acc declare for VLAs

[Tobias Burnus]

LGTM.

Thanks,

Tobias

On 10/6/20 3:28 PM, Tom de Vries wrote:
> Hi,
>
> Consider test-case test.c, with VLA A:
> ...
> int main (void) {
>    int N = 1000;
>    int A[N];
>    #pragma acc declare copy(A)
>    return 0;
> }
> ...
> compiled using:
> ...
> $ gcc test.c -fopenacc -S -fdump-tree-all
> ...
>
> At original, we have:
> ...
>    #pragma acc declare map(tofrom:A);
> ...
> but at gimple, we have a map (to:A.1), but not a map (from:A.1):
> ...
>    int[0:D.2074] * A.1;
>
>    {
>      int A[0:D.2074] [value-expr: *A.1];
>
>      saved_stack.2 = __builtin_stack_save ();
>      try
>        {
>          A.1 = __builtin_alloca_with_align (D.2078, 32);
>          #pragma omp target oacc_declare map(to:(*A.1) [len: D.2076])
>        }
>      finally
>        {
>          __builtin_stack_restore (saved_stack.2);
>        }
>    }
> ...
>
> This is caused by the following incompatibility.  When storing the desired
> from clause in oacc_declare_returns, we use 'A.1' as the key:
> ...
> 10898                 oacc_declare_returns->put (decl, c);
> (gdb) call debug_generic_expr (decl)
> A.1
> (gdb) call debug_generic_expr (c)
> map(from:(*A.1))
> ...
> but when looking it up, we use 'A' as the key:
> ...
> (gdb)
> 1471                  tree *c = oacc_declare_returns->get (t);
> (gdb) call debug_generic_expr (t)
> A
> ...
>
> Fix this by extracing the 'A.1' lookup key from 'A' using the decl-expr.
>
> In addition, unshare the looked up value, to fix avoid running into
> an "incorrect sharing of tree nodes" error.
>
> Using these two fixes, we get our desired:
> ...
>       finally
>         {
> +        #pragma omp target oacc_declare map(from:(*A.1))
>           __builtin_stack_restore (saved_stack.2);
>         }
> ...
>
> Build on x86_64-linux with nvptx accelerator, tested libgomp.
>
> OK for trunk?
>
> Thanks,
> - Tom
>
> [openacc] Fix acc declare for VLAs
>
> gcc/ChangeLog:
>
> 2020-10-06  Tom de Vries  <tdevries@suse.de>
>
>       PR middle-end/90861
>       * gimplify.c (gimplify_bind_expr): Handle lookup in
>       oacc_declare_returns using key with decl-expr.
>
> libgomp/ChangeLog:
>
> 2020-10-06  Tom de Vries  <tdevries@suse.de>
>
>       PR middle-end/90861
>       * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Remove xfail.
>
> ---
>   gcc/gimplify.c                                            | 13 ++++++++++---
>   libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c |  5 -----
>   2 files changed, 10 insertions(+), 8 deletions(-)
>
> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
> index 2dea03cce3d..fa89e797940 100644
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -1468,15 +1468,22 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
>
>         if (flag_openacc && oacc_declare_returns != NULL)
>           {
> -           tree *c = oacc_declare_returns->get (t);
> +           tree key = t;
> +           if (DECL_HAS_VALUE_EXPR_P (key))
> +             {
> +               key = DECL_VALUE_EXPR (key);
> +               if (TREE_CODE (key) == INDIRECT_REF)
> +                 key = TREE_OPERAND (key, 0);
> +             }
> +           tree *c = oacc_declare_returns->get (key);
>             if (c != NULL)
>               {
>                 if (ret_clauses)
>                   OMP_CLAUSE_CHAIN (*c) = ret_clauses;
>
> -               ret_clauses = *c;
> +               ret_clauses = unshare_expr (*c);
>
> -               oacc_declare_returns->remove (t);
> +               oacc_declare_returns->remove (key);
>
>                 if (oacc_declare_returns->is_empty ())
>                   {
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
> index 0f51badca42..714935772c1 100644
> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
> @@ -59,8 +59,3 @@ main ()
>
>     return 0;
>   }
> -
> -
> -/* { dg-xfail-run-if "TODO PR90861" { *-*-* } { "-DACC_MEM_SHARED=0" } }
> -   This might XPASS if the compiler happens to put the two 'A' VLAs at the same
> -   address.  */
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter