| 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; +}
Hi, if we compile the testcase pr84592-2.c from the patch: ... #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; } ... for nvptx offloading, we run into an assert: ... lto1: internal compiler error: in input_varpool_node, at lto-cgraph.c:1424 0x959ebb input_varpool_node gcc/lto-cgraph.c:1422 0x959ebb input_cgraph_1 gcc/lto-cgraph.c:1544 0x959ebb input_symtab() gcc/lto-cgraph.c:1858 0x5aceac read_cgraph_and_symbols gcc/lto/lto.c:2891 0x5aceac lto_main() gcc/lto/lto.c:3356 ... The assert we run into is: ... 1422 gcc_assert (flag_ltrans 1423 || (!node->in_other_partition 1424 && !node->used_from_other_partition)); ... where node is: ... (gdb) call debug_generic_expr (node.decl) test ... and the reason the assert triggers is: ... (gdb) p node.in_other_partition $1 = 1 ... AFAIU, what this means is that the variable test is placed in a different partition than the offloading function main._omp_fn.0 that uses the variable. I looked at where global variables are put into offload_vars, and found that that happens in varpool_node::get_create: ... if ((flag_openacc || flag_openmp) && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) { node->offloadable = 1; if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) { g->have_offload = true; if (!in_lto_p) vec_safe_push (offload_vars, decl); } } ... The patch fixes the ICE there by marking the varpool_node test as offloadable as well. Build and reg-tested libgomp on x86_64 with nvptx accelerator. Bootstrapped and reg-tested on x86_64. OK for stage4 trunk? Thanks, - Tom