| Message ID | 568BC1F5.1010101@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
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. >
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
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.
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
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
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]; } } }
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: