* [PATCH][openacc] Fix acc declare for VLAs
@ 2020-10-06 13:28 Tom de Vries
2020-10-06 14:44 ` Tobias Burnus
2020-10-06 21:44 ` Tobias Burnus
0 siblings, 2 replies; 3+ messages in thread
From: Tom de Vries @ 2020-10-06 13:28 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
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. */
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [PATCH][openacc] Fix acc declare for VLAs
2020-10-06 13:28 [PATCH][openacc] Fix acc declare for VLAs Tom de Vries
@ 2020-10-06 14:44 ` Tobias Burnus
2020-10-06 21:44 ` Tobias Burnus
1 sibling, 0 replies; 3+ messages in thread
From: Tobias Burnus @ 2020-10-06 14:44 UTC (permalink / raw)
To: Tom de Vries, gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
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
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [PATCH][openacc] Fix acc declare for VLAs
2020-10-06 13:28 [PATCH][openacc] Fix acc declare for VLAs Tom de Vries
2020-10-06 14:44 ` Tobias Burnus
@ 2020-10-06 21:44 ` Tobias Burnus
1 sibling, 0 replies; 3+ messages in thread
From: Tobias Burnus @ 2020-10-06 21:44 UTC (permalink / raw)
To: Tom de Vries, gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 878 bytes --]
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
[-- Attachment #2: commit.diff --]
[-- Type: text/x-patch, Size: 1029 bytes --]
commit a9802204603616df14ed47d05f1b86f1bd08d8fb
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Tue Oct 6 23:34:21 2020 +0200
c-c++-common/goacc/declare-pr90861.c: Remove xfail
gcc/testsuite/ChangeLog
PR middle-end/90861
* c-c++-common/goacc/declare-pr90861.c: Remove xfail.
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-pr90861.c b/gcc/testsuite/c-c++-common/goacc/declare-pr90861.c
index 7c905624f7a..c5487bdc8ba 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-pr90861.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-pr90861.c
@@ -17,5 +17,5 @@ void f2 (void)
int A_f2[N_f2];
#pragma acc declare copy(A_f2)
/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_declare map\(to:\(\*A_f2} 1 gimple } }
- { dg-final { scan-tree-dump-times {#pragma omp target oacc_declare map\(from:\(\*A_f2} 1 gimple { xfail *-*-* } } } TODO PR90861 */
+ { dg-final { scan-tree-dump-times {#pragma omp target oacc_declare map\(from:\(\*A_f2} 1 gimple } } */
}
^ permalink raw reply [flat|nested] 3+ messages in thread
end of thread, other threads:[~2020-10-06 21:44 UTC | newest]
Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-10-06 13:28 [PATCH][openacc] Fix acc declare for VLAs Tom de Vries
2020-10-06 14:44 ` Tobias Burnus
2020-10-06 21:44 ` Tobias Burnus
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).