aboutsummaryrefslogtreecommitdiff
path: root/gcc
diff options
context:
space:
mode:
authorTom de Vries <tdevries@suse.de>2020-10-06 13:07:25 +0200
committerTom de Vries <tdevries@suse.de>2020-10-06 16:50:22 +0200
commit3f2e15c2e66af9cca1dfe24ad7e9692f511ebd06 (patch)
treeb4ab82a70f4a4562ff5f74f7c9de4952e6fe81ad /gcc
parent190c04ba36d9c6c3dce41f12012aa97c6d7f22f5 (diff)
downloadgcc-3f2e15c2e66af9cca1dfe24ad7e9692f511ebd06.zip
gcc-3f2e15c2e66af9cca1dfe24ad7e9692f511ebd06.tar.gz
gcc-3f2e15c2e66af9cca1dfe24ad7e9692f511ebd06.tar.bz2
[openacc] Fix acc declare for VLAs
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. 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.
Diffstat (limited to 'gcc')
-rw-r--r--gcc/gimplify.c13
1 files changed, 10 insertions, 3 deletions
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 2dea03cc..fa89e79 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 ())
{