diff options
| author | Jakub Jelinek <jakub@redhat.com> | 2016-01-25 22:34:06 +0100 |
|---|---|---|
| committer | Jakub Jelinek <jakub@gcc.gnu.org> | 2016-01-25 22:34:06 +0100 |
| commit | d902b330dc366702c0f6ab7f529c46cba253427c (patch) | |
| tree | f92373e34a68fd8cf0e9de05f52f0deee1c864be | |
| parent | a4ea172334a511739c98a1abe725363cf3ab096c (diff) | |
| download | gcc-d902b330dc366702c0f6ab7f529c46cba253427c.zip gcc-d902b330dc366702c0f6ab7f529c46cba253427c.tar.gz gcc-d902b330dc366702c0f6ab7f529c46cba253427c.tar.bz2 | |
omp-low.c (lower_omp_target): Set DECL_VALUE_EXPR of new_var even for the non-array case.
* omp-low.c (lower_omp_target) <case USE_DEVICE_PTR>: Set
DECL_VALUE_EXPR of new_var even for the non-array case. Look
through DECL_VALUE_EXPR for expansion.
* c-c++-common/goacc/use_device-1.c: New test.
From-SVN: r232804
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/omp-low.c | 9 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 4 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/use_device-1.c | 14 |
4 files changed, 33 insertions, 0 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 7e63b59..00652e8 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2016-01-25 Jakub Jelinek <jakub@redhat.com> + + * omp-low.c (lower_omp_target) <case USE_DEVICE_PTR>: Set + DECL_VALUE_EXPR of new_var even for the non-array case. Look + through DECL_VALUE_EXPR for expansion. + 2016-01-25 Bernd Edlinger <bernd.edlinger@hotmail.de> * config/mips/mips.c (mips_compute_frame_info): Skip re-computing diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 2977082..3615cbc 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -15878,6 +15878,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) SET_DECL_VALUE_EXPR (new_var, x); DECL_HAS_VALUE_EXPR_P (new_var) = 1; } + else + { + tree new_var = lookup_decl (var, ctx); + x = create_tmp_var_raw (TREE_TYPE (new_var), get_name (new_var)); + gimple_add_tmp_var (x); + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } break; } @@ -16493,6 +16501,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) x = build_fold_addr_expr (v); } } + new_var = DECL_VALUE_EXPR (new_var); x = fold_convert (TREE_TYPE (new_var), x); gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); gimple_seq_add_stmt (&new_body, diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 7e5daa9..5f877b3 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,7 @@ +2016-01-25 Jakub Jelinek <jakub@redhat.com> + + * c-c++-common/goacc/use_device-1.c: New test. + 2016-01-25 Jeff Law <law@redhat.com> PR tree-optimization/69196 diff --git a/gcc/testsuite/c-c++-common/goacc/use_device-1.c b/gcc/testsuite/c-c++-common/goacc/use_device-1.c new file mode 100644 index 0000000..9a4f6d0 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/use_device-1.c @@ -0,0 +1,14 @@ +/* { dg-do compile } */ + +void bar (float *, float *); + +void +foo (float *x, float *y) +{ + int n = 1 << 10; +#pragma acc data create(x[0:n]) copyout(y[0:n]) + { +#pragma acc host_data use_device(x,y) + bar (x, y); + } +} |