[PATCH] OpenACC use_device clause ICE fix

[PATCH] OpenACC use_device clause ICE fix

[Chung-Lin Tang]

[-- Attachment #1: Type: text/plain, Size: 1117 bytes --]

Hi,
we've been encountering an ICE for OpenACC host_data sections, which
has a use_device() clause similar to OpenMP use_device_ptr.

The ICE happens in make_decl_rtl() for scan-created variables, which IIUC,
should not be entered at all for automatic variables.

I believe the problem is, unlike other variable creation cases where the
code is split out into an offloaded child function, a host_data section
is actually host side code, so the child function local variable processing
doesn't apply here; the use_device() referenced variable has to be added
to the current host function.

So here is the quite small fix. This fixed the ICE for OpenACC on trunk
and gomp4. However when I tested it for OpenMP using the case that Julian
provided here[1], the same ICE appeared to be already fixed. I'm not sure
if some other interim change covered it up for OpenMP.

This patch was tested on trunk without regressions. Okay for trunk?

[1] https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00104.html

Thanks,
Chung-Lin

	* omp-low.c (scan_sharing_clauses): Call add_local_decl() for
	use_device/use_device_ptr variables.


[-- Attachment #2: x.diff --]
[-- Type: text/plain, Size: 604 bytes --]

Index: omp-low.c
===================================================================
--- omp-low.c	(revision 232047)
+++ omp-low.c	(working copy)
@@ -1972,7 +1972,10 @@ scan_sharing_clauses (tree clauses, omp_context *c
 	      gcc_assert (DECL_P (decl2));
 	      install_var_local (decl2, ctx);
 	    }
-	  install_var_local (decl, ctx);
+	  decl = install_var_local (decl, ctx);
+	  /* use_device/use_device_ptr items are actually host side variables,
+	     not on the offloaded target; add to current function here.  */
+	  add_local_decl (cfun, decl);
 	  break;
 
 	case OMP_CLAUSE_IS_DEVICE_PTR:

Re: [PATCH] OpenACC use_device clause ICE fix

[Chung-Lin Tang]

Ping.

On 2016/1/5 10:15 PM, Chung-Lin Tang wrote:
> Hi,
> we've been encountering an ICE for OpenACC host_data sections, which
> has a use_device() clause similar to OpenMP use_device_ptr.
> 
> The ICE happens in make_decl_rtl() for scan-created variables, which IIUC,
> should not be entered at all for automatic variables.
> 
> I believe the problem is, unlike other variable creation cases where the
> code is split out into an offloaded child function, a host_data section
> is actually host side code, so the child function local variable processing
> doesn't apply here; the use_device() referenced variable has to be added
> to the current host function.
> 
> So here is the quite small fix. This fixed the ICE for OpenACC on trunk
> and gomp4. However when I tested it for OpenMP using the case that Julian
> provided here[1], the same ICE appeared to be already fixed. I'm not sure
> if some other interim change covered it up for OpenMP.
> 
> This patch was tested on trunk without regressions. Okay for trunk?
> 
> [1] https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00104.html
> 
> Thanks,
> Chung-Lin
> 
> 	* omp-low.c (scan_sharing_clauses): Call add_local_decl() for
> 	use_device/use_device_ptr variables.
> 

Re: [PATCH] OpenACC use_device clause ICE fix

[Bernd Schmidt]

On 01/05/2016 02:15 PM, Chung-Lin Tang wrote:
> 	* omp-low.c (scan_sharing_clauses): Call add_local_decl() for
> 	use_device/use_device_ptr variables.

It looks vaguely plausible, but if everything is part of the host 
function, why make a copy of the decl at all? I.e. what happens if you 
just remove the install_var_local call?


Bernd

Re: [PATCH] OpenACC use_device clause ICE fix

[Chung-Lin Tang]

On 2016/1/20 09:17 PM, Bernd Schmidt wrote:
> On 01/05/2016 02:15 PM, Chung-Lin Tang wrote:
>>     * omp-low.c (scan_sharing_clauses): Call add_local_decl() for
>>     use_device/use_device_ptr variables.
> 
> It looks vaguely plausible, but if everything is part of the host
> function, why make a copy of the decl at all? I.e. what happens if you
> just remove the install_var_local call?

Because (only) inside the OpenMP context, the variable is supposed to
contain the device-side value; a runtime call is used to obtain the
value from the device back to host.  So a new variable is created, the
remap_decl mechanisms are used to change references inside the omp
context, and other references of the original variable are not touched.

Re: [PATCH] OpenACC use_device clause ICE fix

[Bernd Schmidt]

On 01/21/2016 03:22 PM, Chung-Lin Tang wrote:
> On 2016/1/20 09:17 PM, Bernd Schmidt wrote:
>> On 01/05/2016 02:15 PM, Chung-Lin Tang wrote:
>>>      * omp-low.c (scan_sharing_clauses): Call add_local_decl() for
>>>      use_device/use_device_ptr variables.
>>
>> It looks vaguely plausible, but if everything is part of the host
>> function, why make a copy of the decl at all? I.e. what happens if you
>> just remove the install_var_local call?
>
> Because (only) inside the OpenMP context, the variable is supposed to
> contain the device-side value; a runtime call is used to obtain the
> value from the device back to host.  So a new variable is created, the
> remap_decl mechanisms are used to change references inside the omp
> context, and other references of the original variable are not touched.

Hmm, ok. Let's go with your patch then.


Bernd

Re: [PATCH] OpenACC use_device clause ICE fix

[Jakub Jelinek]

On Thu, Jan 21, 2016 at 10:22:19PM +0800, Chung-Lin Tang wrote:
> On 2016/1/20 09:17 PM, Bernd Schmidt wrote:
> > On 01/05/2016 02:15 PM, Chung-Lin Tang wrote:
> >>     * omp-low.c (scan_sharing_clauses): Call add_local_decl() for
> >>     use_device/use_device_ptr variables.
> > 
> > It looks vaguely plausible, but if everything is part of the host
> > function, why make a copy of the decl at all? I.e. what happens if you
> > just remove the install_var_local call?
> 
> Because (only) inside the OpenMP context, the variable is supposed to
> contain the device-side value; a runtime call is used to obtain the
> value from the device back to host.  So a new variable is created, the
> remap_decl mechanisms are used to change references inside the omp
> context, and other references of the original variable are not touched.

The patch looks wrong to me, the var shouldn't be actually used,
it is supposed to have DECL_VALUE_EXPR set for it during omp lowering and
the following gimplification is supposed to replace it.

I've tried the testcases you've listed and couldn't get an ICE, so, if you
see some ICE, can you mail the testcase (in patch form)?
Perhaps there is something wrong with the OpenACC lowering?

	Jakub

Re: [PATCH] OpenACC use_device clause ICE fix

[Chung-Lin Tang]

[-- Attachment #1: Type: text/plain, Size: 1433 bytes --]

On 2016/1/22 12:32 AM, Jakub Jelinek wrote:
> On Thu, Jan 21, 2016 at 10:22:19PM +0800, Chung-Lin Tang wrote:
>> On 2016/1/20 09:17 PM, Bernd Schmidt wrote:
>>> On 01/05/2016 02:15 PM, Chung-Lin Tang wrote:
>>>>     * omp-low.c (scan_sharing_clauses): Call add_local_decl() for
>>>>     use_device/use_device_ptr variables.
>>>
>>> It looks vaguely plausible, but if everything is part of the host
>>> function, why make a copy of the decl at all? I.e. what happens if you
>>> just remove the install_var_local call?
>>
>> Because (only) inside the OpenMP context, the variable is supposed to
>> contain the device-side value; a runtime call is used to obtain the
>> value from the device back to host.  So a new variable is created, the
>> remap_decl mechanisms are used to change references inside the omp
>> context, and other references of the original variable are not touched.
> 
> The patch looks wrong to me, the var shouldn't be actually used,
> it is supposed to have DECL_VALUE_EXPR set for it during omp lowering and
> the following gimplification is supposed to replace it.
> 
> I've tried the testcases you've listed and couldn't get an ICE, so, if you
> see some ICE, can you mail the testcase (in patch form)?
> Perhaps there is something wrong with the OpenACC lowering?
> 
> 	Jakub
> 

I've attached a small testcase that triggers the ICE under -fopenacc. This stll
happens under current trunk.

Thanks,
Chung-Lin


[-- Attachment #2: j.c --]
[-- Type: text/plain, Size: 223 bytes --]


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];
    }
  }
}

Re: [PATCH] OpenACC use_device clause ICE fix

[Jakub Jelinek]

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  <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.

--- 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

Re: [PATCH] OpenACC use_device clause ICE fix

[Jakub Jelinek]

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.

	Jakub

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  <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.

--- 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

Re: [PATCH] OpenACC use_device clause ICE fix

[Chung-Lin Tang]

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