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 <tdevr...@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 <tdevr...@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
