Message ID | 20201006132849.GA16366@delia |
---|---|
State | New |
Headers | show |
Series | [openacc] Fix acc declare for VLAs | expand |
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
And as spotted by Thomas, Tom's patch also resolved an XFAIL in gcc/testsuite. Committed as r11-3687-ga9802204603616df14ed47d05f1b86f1bd08d8fb after testing it on x86-64-gnu-linux. Tobias On 10/6/20 3:28 PM, Tom de Vries wrote: ... > [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. ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
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. */