Re: [PATCH] OpenACC use_device clause ICE fix

[Chung-Lin Tang]

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


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