diff mbox

OpenACC use_device clause ICE fix

Message ID 20160125100647.GI3017@tucnak.redhat.com
State New
Headers show

Commit Message

Jakub Jelinek Jan. 25, 2016, 10:06 a.m. UTC 
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  <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.

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



	Jakub

Comments

Chung-Lin Tang Jan. 26, 2016, 5:39 a.m. UTC | #1 
On 2016/1/25 7:06 PM, Jakub Jelinek wrote:
> The following ICEs without the patch and works with it, so I think it is
> better:
> 
> 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.
> 
> 	* c-c++-common/goacc/use_device-1.c: New test.

Thanks, the test was indeed just a reduction of a whole example program, which I'm not sure
we're at liberty to directly include in the testsuite. I've verified that the patch
allows the program to build and run correctly.

Thanks,
Chung-Lin
diff mbox

Patch

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