Re: [PATCH] OpenACC use_device clause ICE fix

[Jakub Jelinek]

On Mon, Jan 25, 2016 at 11:02:05AM +0100, Jakub Jelinek wrote:
> On Mon, Jan 25, 2016 at 10:58:17AM +0100, Jakub Jelinek wrote:
> > --- 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];
> > +    }
> > +  }
> > +}
> 
> Though the testcase looks invalid to me, how can you dereference
> the device pointer on the host?  Though, for a testcase that it doesn't ICE
> maybe good enough.

The following ICEs without the patch and works with it, so I think it is
better:

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,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);
+  }
+}


        Jakub