omp-low.c (lower_omp_target): Set DECL_VALUE_EXPR of new_var even for the non-array...

	* 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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 7e63b598e00047fd4a21bf95aea2987a919c9e42..00652e83923c628b361da48f0e7bcf07521817d5 100644 (file)
--- 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 29770820a06b6f73147582868b809ebac9829410..3615cbc5f9e7efdceb8fc515ba599f1fd2014df8 100644 (file)
--- 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 7e5daa9f3d4629fc357567e1e0f9f789c414f29a..5f877b376724aa298adb00767a91e8938aa1e631 100644 (file)
--- 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 (file)
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);
+  }
+}