On Mon, Jan 25, 2016 at 05:52:56PM +0900, Chung-Lin Tang wrote:
> I've attached a small testcase that triggers the ICE under -fopenacc. This 
> stll
> happens under current trunk.

Then I think I'd prefer (untested so far):

2016-01-25  Jakub Jelinek  <ja...@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.

        * c-c++-common/goacc/use_device-1.c: New test.

--- gcc/omp-low.c.jj    2016-01-21 00:55:19.000000000 +0100
+++ gcc/omp-low.c       2016-01-25 10:45:30.995510057 +0100
@@ -15878,6 +15878,14 @@ lower_omp_target (gimple_stmt_iterator *
            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 *
                        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,
--- gcc/testsuite/c-c++-common/goacc/use_device-1.c.jj  2016-01-25 
10:56:33.472310437 +0100
+++ gcc/testsuite/c-c++-common/goacc/use_device-1.c     2016-01-25 
10:56:43.128176481 +0100
@@ -0,0 +1,15 @@
+/* { dg-do compile } */
+
+void
+foo (float *x, float *y)
+{
+  int n = 1 << 20;
+#pragma acc data create(x[0:n]) copyout(y[0:n])
+  {
+#pragma acc host_data use_device(x,y)
+    {
+      for (int i = 1; i < n; i++)
+       y[0] += x[i] * y[i];
+    }
+  }
+}

        Jakub

Reply via email to