Message ID | 2299637a-3124-8609-e68d-225703f16abe@mentor.com |
---|---|
State | New |
Headers | show |
Series | Fix ICE for static vars in offloaded functions | expand |
On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote: > Fix ICE for static vars in offloaded functions > > 2018-03-06 Tom de Vries <tom@codesourcery.com> > > PR lto/84592 > * varpool.c (varpool_node::get_create): Mark static variables in > offloaded functions as offloadable. > > * testsuite/libgomp.c/pr84592-2.c: New test. > * testsuite/libgomp.c/pr84592.c: New test. > * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. Ok, thanks Jakub
On Wed, 7 Mar 2018, Jakub Jelinek wrote: > On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote: > > Fix ICE for static vars in offloaded functions > > > > 2018-03-06 Tom de Vries <tom@codesourcery.com> > > > > PR lto/84592 > > * varpool.c (varpool_node::get_create): Mark static variables in > > offloaded functions as offloadable. > > > > * testsuite/libgomp.c/pr84592-2.c: New test. > > * testsuite/libgomp.c/pr84592.c: New test. > > * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. > > Ok, thanks + bool in_offload_func + = (cfun + && TREE_STATIC (decl) + && (lookup_attribute ("omp target entr I think you want to use decl_function_context (decl) here, not rely on magic cfun being set. The whole varpool.c file doesn't mention cfun yet and you shoudln't either. please fix if you already committed the fix. Thanks, Richard. > Jakub > >
On Wed, Mar 07, 2018 at 02:29:48PM +0100, Richard Biener wrote: > On Wed, 7 Mar 2018, Jakub Jelinek wrote: > > > On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote: > > > Fix ICE for static vars in offloaded functions > > > > > > 2018-03-06 Tom de Vries <tom@codesourcery.com> > > > > > > PR lto/84592 > > > * varpool.c (varpool_node::get_create): Mark static variables in > > > offloaded functions as offloadable. > > > > > > * testsuite/libgomp.c/pr84592-2.c: New test. > > > * testsuite/libgomp.c/pr84592.c: New test. > > > * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. > > > > Ok, thanks > > + bool in_offload_func > + = (cfun > + && TREE_STATIC (decl) > + && (lookup_attribute ("omp target entr > > I think you want to use decl_function_context (decl) here, > not rely on magic cfun being set. The whole varpool.c file > doesn't mention cfun yet and you shoudln't either. Oops, sure, thanks for catching it. Jakub
On 03/07/2018 02:29 PM, Richard Biener wrote: > On Wed, 7 Mar 2018, Jakub Jelinek wrote: > >> On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote: >>> Fix ICE for static vars in offloaded functions >>> >>> 2018-03-06 Tom de Vries <tom@codesourcery.com> >>> >>> PR lto/84592 >>> * varpool.c (varpool_node::get_create): Mark static variables in >>> offloaded functions as offloadable. >>> >>> * testsuite/libgomp.c/pr84592-2.c: New test. >>> * testsuite/libgomp.c/pr84592.c: New test. >>> * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. >> >> Ok, thanks > > + bool in_offload_func > + = (cfun > + && TREE_STATIC (decl) > + && (lookup_attribute ("omp target entr > > I think you want to use decl_function_context (decl) here, > not rely on magic cfun being set. The whole varpool.c file > doesn't mention cfun yet and you shoudln't either. > decl_function_context (decl) returns main: ... (gdb) call debug_generic_expr (decl) test (gdb) call decl_function_context (decl) $2 = (tree_node *) 0x7ffff6978c00 (gdb) call debug_generic_expr ($2) main ... while the function annotated as being an offload function is main._omp_fn.0. The varpool_node::get_create is called during cgraph_edge::rebuild_edges here in expand_omp_target: ... 7087 /* Fix the callgraph edges for child_cfun. Those for cfun will be 7088 fixed in a following pass. */ 7089 push_cfun (child_cfun); 7090 if (need_asm) 7091 assign_assembler_name_if_needed (child_fn); 7092 cgraph_edge::rebuild_edges (); ... Thanks, - Tom
On Wed, 7 Mar 2018, Tom de Vries wrote: > On 03/07/2018 02:29 PM, Richard Biener wrote: > > On Wed, 7 Mar 2018, Jakub Jelinek wrote: > > > > > On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote: > > > > Fix ICE for static vars in offloaded functions > > > > > > > > 2018-03-06 Tom de Vries <tom@codesourcery.com> > > > > > > > > PR lto/84592 > > > > * varpool.c (varpool_node::get_create): Mark static variables in > > > > offloaded functions as offloadable. > > > > > > > > * testsuite/libgomp.c/pr84592-2.c: New test. > > > > * testsuite/libgomp.c/pr84592.c: New test. > > > > * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. > > > > > > Ok, thanks > > > > + bool in_offload_func > > + = (cfun > > + && TREE_STATIC (decl) > > + && (lookup_attribute ("omp target entr > > > > I think you want to use decl_function_context (decl) here, > > not rely on magic cfun being set. The whole varpool.c file > > doesn't mention cfun yet and you shoudln't either. > > > > decl_function_context (decl) returns main: > ... > (gdb) call debug_generic_expr (decl) > test > (gdb) call decl_function_context (decl) > $2 = (tree_node *) 0x7ffff6978c00 > (gdb) call debug_generic_expr ($2) > main > ... > while the function annotated as being an offload function is main._omp_fn.0. Well, that's because the static isn't duplicated (it can't be) so it retains the original context. > The varpool_node::get_create is called during cgraph_edge::rebuild_edges here > in expand_omp_target: But at this point it's not created but just looked up, right? I think the fix is to mark the decl as offloaded when we walk the IL of the outlined function. The current point looks like a hack. Richard. > ... > 7087 /* Fix the callgraph edges for child_cfun. Those for cfun will > be > 7088 fixed in a following pass. */ > 7089 push_cfun (child_cfun); > 7090 if (need_asm) > 7091 assign_assembler_name_if_needed (child_fn); > 7092 cgraph_edge::rebuild_edges (); > ... > > Thanks, > - Tom > >
On 03/07/2018 04:01 PM, Richard Biener wrote: > On Wed, 7 Mar 2018, Tom de Vries wrote: > >> On 03/07/2018 02:29 PM, Richard Biener wrote: >>> On Wed, 7 Mar 2018, Jakub Jelinek wrote: >>> >>>> On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote: >>>>> Fix ICE for static vars in offloaded functions >>>>> >>>>> 2018-03-06 Tom de Vries <tom@codesourcery.com> >>>>> >>>>> PR lto/84592 >>>>> * varpool.c (varpool_node::get_create): Mark static variables in >>>>> offloaded functions as offloadable. >>>>> >>>>> * testsuite/libgomp.c/pr84592-2.c: New test. >>>>> * testsuite/libgomp.c/pr84592.c: New test. >>>>> * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. >>>> >>>> Ok, thanks >>> >>> + bool in_offload_func >>> + = (cfun >>> + && TREE_STATIC (decl) >>> + && (lookup_attribute ("omp target entr >>> >>> I think you want to use decl_function_context (decl) here, >>> not rely on magic cfun being set. The whole varpool.c file >>> doesn't mention cfun yet and you shoudln't either. >>> >> >> decl_function_context (decl) returns main: >> ... >> (gdb) call debug_generic_expr (decl) >> test >> (gdb) call decl_function_context (decl) >> $2 = (tree_node *) 0x7ffff6978c00 >> (gdb) call debug_generic_expr ($2) >> main >> ... >> while the function annotated as being an offload function is main._omp_fn.0. > > Well, that's because the static isn't duplicated (it can't be) so it > retains the original context. > [ Actually the static is duplicated in replace_by_duplicate_decl, but the statements using it are not rewritten to use the duplicate, so indeed, effectively it's not duplicated. ] >> The varpool_node::get_create is called during cgraph_edge::rebuild_edges here >> in expand_omp_target: > > But at this point it's not created but just looked up, right? > No, the varpool_node is created at that point. > I think the fix is to mark the decl as offloaded when we walk the IL > of the outlined function. The current point looks like a hack. > OK, I'll try to find a better fix location. Thanks, - Tom > Richard. > >> ... >> 7087 /* Fix the callgraph edges for child_cfun. Those for cfun will >> be >> 7088 fixed in a following pass. */ >> 7089 push_cfun (child_cfun); >> 7090 if (need_asm) >> 7091 assign_assembler_name_if_needed (child_fn); >> 7092 cgraph_edge::rebuild_edges (); >> ... >> >> Thanks, >> - Tom >> >> >
Fix ICE for static vars in offloaded functions 2018-03-06 Tom de Vries <tom@codesourcery.com> PR lto/84592 * varpool.c (varpool_node::get_create): Mark static variables in offloaded functions as offloadable. * testsuite/libgomp.c/pr84592-2.c: New test. * testsuite/libgomp.c/pr84592.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. --- gcc/varpool.c | 18 +++++++++--- libgomp/testsuite/libgomp.c/pr84592-2.c | 20 ++++++++++++++ libgomp/testsuite/libgomp.c/pr84592.c | 32 ++++++++++++++++++++++ .../libgomp.oacc-c-c++-common/pr84592-3.c | 32 ++++++++++++++++++++++ 4 files changed, 98 insertions(+), 4 deletions(-) diff --git a/gcc/varpool.c b/gcc/varpool.c index 418753cca2a..a4fd892ca4d 100644 --- a/gcc/varpool.c +++ b/gcc/varpool.c @@ -151,11 +151,21 @@ varpool_node::get_create (tree decl) node = varpool_node::create_empty (); node->decl = decl; - if ((flag_openacc || flag_openmp) - && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + if (flag_openacc || flag_openmp) { - node->offloadable = 1; - if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) + bool offload_var + = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)); + bool in_offload_func + = (cfun + && TREE_STATIC (decl) + && (lookup_attribute ("omp target entrypoint", + DECL_ATTRIBUTES (cfun->decl)) + || lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (cfun->decl)))); + if (offload_var || in_offload_func) + node->offloadable = 1; + + if (offload_var && ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) { g->have_offload = true; if (!in_lto_p) diff --git a/libgomp/testsuite/libgomp.c/pr84592-2.c b/libgomp/testsuite/libgomp.c/pr84592-2.c new file mode 100644 index 00000000000..021497b28ff --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr84592-2.c @@ -0,0 +1,20 @@ +#include <stdlib.h> + +int +main (void) +{ + int n[1]; + + n[0] = 3; + +#pragma omp target + { + static int test[4] = { 1, 2, 3, 4 }; + n[0] += test[n[0]]; + } + + if (n[0] != 7) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr84592.c b/libgomp/testsuite/libgomp.c/pr84592.c new file mode 100644 index 00000000000..197fd19bacc --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr84592.c @@ -0,0 +1,32 @@ +/* { dg-additional-options "-ftree-switch-conversion" } */ + +#include <stdlib.h> + +int +main (void) +{ + int n[1]; + + n[0] = 4; + +#pragma omp target + { + int a = n[0]; + + switch (a & 3) + { + case 0: a = 4; break; + case 1: a = 3; break; + case 2: a = 2; break; + default: + a = 1; break; + } + + n[0] = a; + } + + if (n[0] != 4) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c new file mode 100644 index 00000000000..afcc1de7635 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c @@ -0,0 +1,32 @@ +/* { dg-additional-options "-ftree-switch-conversion" } */ + +#include <stdlib.h> + +#pragma acc routine seq +static int __attribute__((noinline)) foo (int n) +{ + switch (n & 3) + { + case 0: return 4; + case 1: return 3; + case 2: return 2; + default: + return 1; + } +} + +int +main (void) +{ + int n[1]; + n[0] = 4; +#pragma acc parallel copy(n) + { + n[0] = foo (n[0]); + } + + if (n[0] != 4) + abort (); + + return 0; +}