| Message ID | 319b3ebd-c601-449b-718c-963b68414224@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [openacc] Fix PR71959: lto dump of callee counts | expand |
Hi, I have noticed a few things... On Thu, Sep 20 2018, Cesar Philippidis wrote: > This is another old gomp4 patch that demotes an ICE in PR71959 to a > linker warning. One problem here is that it is not clear if OpenACC > allows individual member functions in C++ classes to be marked as acc > routines. There's another issue accessing member data inside offloaded > regions. We'll add some support for member data OpenACC 2.6, but some of > the OpenACC C++ semantics are still unclear. > > Is this OK for trunk? I bootstrapped and regtested it for x86_64 Linux > with nvptx offloading. > > Thanks, > Cesar > [PR71959] lto dump of callee counts > > 2018-XX-YY Nathan Sidwell <nathan@acm.org> > Cesar Philippidis <cesar@codesourcery.com> > > gcc/ > * ipa-inline-analysis.c (inline_write_summary): Only dump callee > counts when dumping the function's body. ...the changelog needs updating and.... > > libgomp/ > * testsuite/libgomp.oacc-c++/pr71959.C: New. > * testsuite/libgomp.oacc-c++/pr71959-a.C: New. > > (cherry picked from gomp-4_0-branch r239788) > --- > gcc/ipa-fnsummary.c | 18 ++++++++--- > .../testsuite/libgomp.oacc-c++/pr71959-a.C | 31 +++++++++++++++++++ > libgomp/testsuite/libgomp.oacc-c++/pr71959.C | 31 +++++++++++++++++++ > 3 files changed, 75 insertions(+), 5 deletions(-) > create mode 100644 libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C > create mode 100644 libgomp/testsuite/libgomp.oacc-c++/pr71959.C > > diff --git a/gcc/ipa-fnsummary.c b/gcc/ipa-fnsummary.c > index 62095c6cf6f..e796b085e14 100644 > --- a/gcc/ipa-fnsummary.c > +++ b/gcc/ipa-fnsummary.c > @@ -3409,8 +3409,10 @@ ipa_fn_summary_write (void) > int i; > size_time_entry *e; > struct condition *c; > + int index = lto_symtab_encoder_encode (encoder, cnode); > + bool body = encoder->nodes[index].body; > > - streamer_write_uhwi (ob, lto_symtab_encoder_encode (encoder, cnode)); > + streamer_write_uhwi (ob, index); > streamer_write_hwi (ob, info->estimated_self_stack_size); > streamer_write_hwi (ob, info->self_size); > info->time.stream_out (ob); > @@ -3453,10 +3455,16 @@ ipa_fn_summary_write (void) > info->array_index->stream_out (ob); > else > streamer_write_uhwi (ob, 0); > - for (edge = cnode->callees; edge; edge = edge->next_callee) > - write_ipa_call_summary (ob, edge); > - for (edge = cnode->indirect_calls; edge; edge = edge->next_callee) > - write_ipa_call_summary (ob, edge); > + if (body) > + { > + /* Only write callee counts when we're emitting the > + body, as the reader only knows about the callees when > + the body's emitted. */ this comment is wrong because write_ipa_call_summary does not only stream counts but also sizes, predicates and other stuff. I also wonder 1) whether the cnode->definition test that guards this streaming should not be replaced by your body check (and why the definition is not enough, really) and 2) why you don't have the same problem with ipa_prop_write_jump_functions and the function it calls. Martin > + for (edge = cnode->callees; edge; edge = edge->next_callee) > + write_ipa_call_summary (ob, edge); > + for (edge = cnode->indirect_calls; edge; edge = edge->next_callee) > + write_ipa_call_summary (ob, edge); > + } > } > } > streamer_write_char_stream (ob->main_stream, 0);
On Tue, 25 Sep 2018 14:59:18 +0200 Martin Jambor <mjambor@suse.cz> wrote: > Hi, > > I have noticed a few things... > > On Thu, Sep 20 2018, Cesar Philippidis wrote: > > This is another old gomp4 patch that demotes an ICE in PR71959 to a > > linker warning. One problem here is that it is not clear if OpenACC > > allows individual member functions in C++ classes to be marked as > > acc routines. There's another issue accessing member data inside > > offloaded regions. We'll add some support for member data OpenACC > > 2.6, but some of the OpenACC C++ semantics are still unclear. > > > > Is this OK for trunk? I bootstrapped and regtested it for x86_64 > > Linux with nvptx offloading. > [...] The testcase associated with this bug appears to be fixed by the following patch: https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01167.html So, it's unclear if there's anything left to do here, and this patch can probably be withdrawn. Thanks, Julian
On Fri, 21 Dec 2018 02:56:36 +0000 Julian Brown <julian@codesourcery.com> wrote: > On Tue, 25 Sep 2018 14:59:18 +0200 > Martin Jambor <mjambor@suse.cz> wrote: > > > Hi, > > > > I have noticed a few things... > > > > On Thu, Sep 20 2018, Cesar Philippidis wrote: > > > This is another old gomp4 patch that demotes an ICE in PR71959 to > > > a linker warning. One problem here is that it is not clear if > > > OpenACC allows individual member functions in C++ classes to be > > > marked as acc routines. There's another issue accessing member > > > data inside offloaded regions. We'll add some support for member > > > data OpenACC 2.6, but some of the OpenACC C++ semantics are still > > > unclear. > > > > > > Is this OK for trunk? I bootstrapped and regtested it for x86_64 > > > Linux with nvptx offloading. > > [...] > > The testcase associated with this bug appears to be fixed by the > following patch: > > https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01167.html > > So, it's unclear if there's anything left to do here, and this patch > can probably be withdrawn. ...or actually, maybe we should keep the new testcase in case of future regressions. This patch contains just that. OK to apply? Thanks, Julian ChangeLog 2018-xx-yy Nathan Sidwell <nathan@acm.org> PR lto/71959 libgomp/ * testsuite/libgomp.oacc-c++/pr71959-a.C: New. * testsuite/libgomp.oacc-c++/pr71959.C: New. commit c69dce8ba0ecd7ff620f4f1b8dacc94c61984107 Author: Julian Brown <julian@codesourcery.com> Date: Wed Dec 19 05:01:58 2018 -0800 Add testcase from PR71959 libgomp/ * testsuite/libgomp.oacc-c++/pr71959-a.C: New. * testsuite/libgomp.oacc-c++/pr71959.C: New. diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C new file mode 100644 index 0000000..ec4b14a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C @@ -0,0 +1,31 @@ +// { dg-do compile } + +struct Iter +{ + int *cursor; + + void ctor (int *cursor_) asm("_ZN4IterC1EPi"); + int *point () const asm("_ZNK4Iter5pointEv"); +}; + +#pragma acc routine +void Iter::ctor (int *cursor_) +{ + cursor = cursor_; +} + +#pragma acc routine +int *Iter::point () const +{ + return cursor; +} + +void apply (int (*fn)(), Iter out) asm ("_ZN5Apply5applyEPFivE4Iter"); + +#pragma acc routine +void apply (int (*fn)(), struct Iter out) +{ *out.point() = fn (); } + +extern "C" void __gxx_personality_v0 () +{ +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959.C b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C new file mode 100644 index 0000000..8508c17 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C @@ -0,0 +1,31 @@ +// { dg-additional-sources "pr71959-a.C" } + +// pr lto/71959 ICEd LTO due to mismatch between writing & reading behaviour + +struct Iter +{ + int *cursor; + + Iter(int *cursor_) : cursor(cursor_) {} + + int *point() const { return cursor; } +}; + +#pragma acc routine seq +int one () { return 1; } + +struct Apply +{ + static void apply (int (*fn)(), Iter out) + { *out.point() = fn (); } +}; + +int main () +{ + int x; + +#pragma acc parallel copyout(x) + Apply::apply (one, Iter (&x)); + + return x != 1; +}
On Fri, Dec 21, 2018 at 01:23:03PM +0000, Julian Brown wrote: > 2018-xx-yy Nathan Sidwell <nathan@acm.org> > > PR lto/71959 > libgomp/ > * testsuite/libgomp.oacc-c++/pr71959-a.C: New. > * testsuite/libgomp.oacc-c++/pr71959.C: New. Just nits, better use pr71959-aux.cc (*.cc files aren't considered as testcases by *.exp: set tests [lsort [concat \ [find $srcdir/$subdir *.C] \ [find $srcdir/$subdir/../libgomp.oacc-c-c++-common *.c]]] ) and just a is weird. > commit c69dce8ba0ecd7ff620f4f1b8dacc94c61984107 > Author: Julian Brown <julian@codesourcery.com> > Date: Wed Dec 19 05:01:58 2018 -0800 > > Add testcase from PR71959 > > libgomp/ Please mention PR lto/71959 here in the ChangeLog. > * testsuite/libgomp.oacc-c++/pr71959-a.C: New. > * testsuite/libgomp.oacc-c++/pr71959.C: New. > +void apply (int (*fn)(), Iter out) asm ("_ZN5Apply5applyEPFivE4Iter"); Will this work even on targets that use _ or other symbol prefixes? > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C > @@ -0,0 +1,31 @@ > +// { dg-additional-sources "pr71959-a.C" } > + > +// pr lto/71959 ICEd LTO due to mismatch between writing & reading behaviour Capital PR instead of pr . Jakub
Hi Jakub, Thanks for review! On Fri, 21 Dec 2018 14:31:19 +0100 Jakub Jelinek <jakub@redhat.com> wrote: > On Fri, Dec 21, 2018 at 01:23:03PM +0000, Julian Brown wrote: > > 2018-xx-yy Nathan Sidwell <nathan@acm.org> > > > > PR lto/71959 > > libgomp/ > > * testsuite/libgomp.oacc-c++/pr71959-a.C: New. > > * testsuite/libgomp.oacc-c++/pr71959.C: New. > > Just nits, better use pr71959-aux.cc (*.cc files aren't considered as > testcases by *.exp: > set tests [lsort [concat \ > [find $srcdir/$subdir *.C] \ > [find > $srcdir/$subdir/../libgomp.oacc-c-c++-common *.c]]] ) and just a is > weird. Fixed. > > commit c69dce8ba0ecd7ff620f4f1b8dacc94c61984107 > > Author: Julian Brown <julian@codesourcery.com> > > Date: Wed Dec 19 05:01:58 2018 -0800 > > > > Add testcase from PR71959 > > > > libgomp/ > > Please mention > PR lto/71959 > here in the ChangeLog. Fixed. > > * testsuite/libgomp.oacc-c++/pr71959-a.C: New. > > * testsuite/libgomp.oacc-c++/pr71959.C: New. > > > +void apply (int (*fn)(), Iter out) asm > > ("_ZN5Apply5applyEPFivE4Iter"); > > Will this work even on targets that use _ or other symbol prefixes? I'd guess so, else there would be no portable way of using "asm" to write pre-mangled C++ names. The only existing similar uses I could find in the testsuite are for the ifunc attribute, not asm, though (e.g. g++.dg/ext/attr-ifunc-*.C). Anyway, OpenACC is only useful for a handful of targets at present, neither of which use special symbol prefixes AFAIK. > > --- /dev/null > > +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C > > @@ -0,0 +1,31 @@ > > +// { dg-additional-sources "pr71959-a.C" } > > + > > +// pr lto/71959 ICEd LTO due to mismatch between writing & reading > > behaviour > > Capital PR instead of pr . Fixed. OK now? Thanks, Julian
Hi Julian, > On 21 Dec 2018, at 16:47, Julian Brown <julian@codesourcery.com> wrote: > > On Fri, 21 Dec 2018 14:31:19 +0100 > Jakub Jelinek <jakub@redhat.com> wrote: > >> On Fri, Dec 21, 2018 at 01:23:03PM +0000, Julian Brown wrote: >>> 2018-xx-yy Nathan Sidwell <nathan@acm.org> > >>> * testsuite/libgomp.oacc-c++/pr71959-a.C: New. >>> * testsuite/libgomp.oacc-c++/pr71959.C: New. >> >>> +void apply (int (*fn)(), Iter out) asm >>> ("_ZN5Apply5applyEPFivE4Iter"); >> >> Will this work even on targets that use _ or other symbol prefixes? > > I'd guess so, else there would be no portable way of using "asm" to > write pre-mangled C++ names. The only existing similar uses I could find > in the testsuite are for the ifunc attribute, not asm, though (e.g. > g++.dg/ext/attr-ifunc-*.C). It won’t work on such targets (e.g. Darwin) … but it’s not too hard to make it happen (see, for example, gcc.dg/memcmp-1.c) One just has to remember that __USER_LABEL_PREFIX__ is a token, not a string. so .. in the example above… #define STR1(X) #X #define STR2(X) STR1(X) …. asm(STR2(__USER_LABEL_PREFIX__) "_ZN5Apply5applyEPFivE4Iter”); > Anyway, OpenACC is only useful for a handful of targets at present, > neither of which use special symbol prefixes AFAIK. I have hopes of one day getting offloading to work on Darwin (the only limitation is developer time, not technical feasibility) .. cheers Iain
On Sat, 22 Dec 2018 15:09:34 +0000 Iain Sandoe <idsandoe@googlemail.com> wrote: > Hi Julian, > > > On 21 Dec 2018, at 16:47, Julian Brown <julian@codesourcery.com> > > wrote: > > > On Fri, 21 Dec 2018 14:31:19 +0100 > > Jakub Jelinek <jakub@redhat.com> wrote: > > > >> On Fri, Dec 21, 2018 at 01:23:03PM +0000, Julian Brown wrote: > >>> 2018-xx-yy Nathan Sidwell <nathan@acm.org> > > > > >>> * testsuite/libgomp.oacc-c++/pr71959-a.C: New. > >>> * testsuite/libgomp.oacc-c++/pr71959.C: New. > >> > >>> +void apply (int (*fn)(), Iter out) asm > >>> ("_ZN5Apply5applyEPFivE4Iter"); > >> > >> Will this work even on targets that use _ or other symbol > >> prefixes? > > > > I'd guess so, else there would be no portable way of using "asm" to > > write pre-mangled C++ names. The only existing similar uses I could > > find in the testsuite are for the ifunc attribute, not asm, though > > (e.g. g++.dg/ext/attr-ifunc-*.C). > > It won’t work on such targets (e.g. Darwin) > … but it’s not too hard to make it happen (see, for example, > gcc.dg/memcmp-1.c) > > One just has to remember that __USER_LABEL_PREFIX__ is a token, not a > string. > > so .. in the example above… > > #define STR1(X) #X > #define STR2(X) STR1(X) > > …. > > asm(STR2(__USER_LABEL_PREFIX__) "_ZN5Apply5applyEPFivE4Iter”); Thanks! I've amended the test to use this technique (though I can't easily test on Darwin, so this is "best effort"). > > Anyway, OpenACC is only useful for a handful of targets at present, > > neither of which use special symbol prefixes AFAIK. > > I have hopes of one day getting offloading to work on Darwin (the > only limitation is developer time, not technical feasibility) .. Is this OK now (for stage 4)? Thanks, Julian commit 2ee3f8d09a7b2af6c9ba29cdd8e8587db1946c0b Author: Julian Brown <julian@codesourcery.com> Date: Wed Dec 19 05:01:58 2018 -0800 Add testcase from PR71959 libgomp/ PR lto/71959 * testsuite/libgomp.oacc-c++/pr71959-aux.cc: New. * testsuite/libgomp.oacc-c++/pr71959.C: New. diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959-aux.cc b/libgomp/testsuite/libgomp.oacc-c++/pr71959-aux.cc new file mode 100644 index 0000000..10a6eeb --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959-aux.cc @@ -0,0 +1,35 @@ +// { dg-do compile } + +#define STR1(X) #X +#define STR2(X) STR1(X) +#define LABEL(X) STR2(__USER_LABEL_PREFIX__) X + +struct Iter +{ + int *cursor; + + void ctor (int *cursor_) asm (LABEL ("_ZN4IterC1EPi")); + int *point () const asm (LABEL ("_ZNK4Iter5pointEv")); +}; + +#pragma acc routine +void Iter::ctor (int *cursor_) +{ + cursor = cursor_; +} + +#pragma acc routine +int *Iter::point () const +{ + return cursor; +} + +void apply (int (*fn)(), Iter out) asm (LABEL ("_ZN5Apply5applyEPFivE4Iter")); + +#pragma acc routine +void apply (int (*fn)(), struct Iter out) +{ *out.point() = fn (); } + +extern "C" void __gxx_personality_v0 () +{ +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959.C b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C new file mode 100644 index 0000000..bf27a75 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C @@ -0,0 +1,31 @@ +// { dg-additional-sources "pr71959-aux.cc" } + +// PR lto/71959 ICEd LTO due to mismatch between writing & reading behaviour + +struct Iter +{ + int *cursor; + + Iter(int *cursor_) : cursor(cursor_) {} + + int *point() const { return cursor; } +}; + +#pragma acc routine seq +int one () { return 1; } + +struct Apply +{ + static void apply (int (*fn)(), Iter out) + { *out.point() = fn (); } +}; + +int main () +{ + int x; + +#pragma acc parallel copyout(x) + Apply::apply (one, Iter (&x)); + + return x != 1; +}
On Tue, Jan 08, 2019 at 01:30:45PM +0000, Julian Brown wrote: > commit 2ee3f8d09a7b2af6c9ba29cdd8e8587db1946c0b > Author: Julian Brown <julian@codesourcery.com> > Date: Wed Dec 19 05:01:58 2018 -0800 > > Add testcase from PR71959 > > libgomp/ > > PR lto/71959 > * testsuite/libgomp.oacc-c++/pr71959-aux.cc: New. > * testsuite/libgomp.oacc-c++/pr71959.C: New. Ok, thanks. Jakub
[PR71959] lto dump of callee counts 2018-XX-YY Nathan Sidwell <nathan@acm.org> Cesar Philippidis <cesar@codesourcery.com> gcc/ * ipa-inline-analysis.c (inline_write_summary): Only dump callee counts when dumping the function's body. libgomp/ * testsuite/libgomp.oacc-c++/pr71959.C: New. * testsuite/libgomp.oacc-c++/pr71959-a.C: New. (cherry picked from gomp-4_0-branch r239788) --- gcc/ipa-fnsummary.c | 18 ++++++++--- .../testsuite/libgomp.oacc-c++/pr71959-a.C | 31 +++++++++++++++++++ libgomp/testsuite/libgomp.oacc-c++/pr71959.C | 31 +++++++++++++++++++ 3 files changed, 75 insertions(+), 5 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C create mode 100644 libgomp/testsuite/libgomp.oacc-c++/pr71959.C diff --git a/gcc/ipa-fnsummary.c b/gcc/ipa-fnsummary.c index 62095c6cf6f..e796b085e14 100644 --- a/gcc/ipa-fnsummary.c +++ b/gcc/ipa-fnsummary.c @@ -3409,8 +3409,10 @@ ipa_fn_summary_write (void) int i; size_time_entry *e; struct condition *c; + int index = lto_symtab_encoder_encode (encoder, cnode); + bool body = encoder->nodes[index].body; - streamer_write_uhwi (ob, lto_symtab_encoder_encode (encoder, cnode)); + streamer_write_uhwi (ob, index); streamer_write_hwi (ob, info->estimated_self_stack_size); streamer_write_hwi (ob, info->self_size); info->time.stream_out (ob); @@ -3453,10 +3455,16 @@ ipa_fn_summary_write (void) info->array_index->stream_out (ob); else streamer_write_uhwi (ob, 0); - for (edge = cnode->callees; edge; edge = edge->next_callee) - write_ipa_call_summary (ob, edge); - for (edge = cnode->indirect_calls; edge; edge = edge->next_callee) - write_ipa_call_summary (ob, edge); + if (body) + { + /* Only write callee counts when we're emitting the + body, as the reader only knows about the callees when + the body's emitted. */ + for (edge = cnode->callees; edge; edge = edge->next_callee) + write_ipa_call_summary (ob, edge); + for (edge = cnode->indirect_calls; edge; edge = edge->next_callee) + write_ipa_call_summary (ob, edge); + } } } streamer_write_char_stream (ob->main_stream, 0); diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C new file mode 100644 index 00000000000..9486512d0e7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C @@ -0,0 +1,31 @@ +// { dg-do compile } + +struct Iter +{ + int *cursor; + + void ctor (int *cursor_) asm("_ZN4IterC1EPi"); + int *point () const asm("_ZNK4Iter5pointEv"); +}; + +#pragma acc routine +void Iter::ctor (int *cursor_) +{ + cursor = cursor_; +} + +#pragma acc routine +int *Iter::point () const +{ + return cursor; +} + +void apply (int (*fn)(), Iter out) asm ("_ZN5Apply5applyEPFivE4Iter"); + +#pragma acc routine +void apply (int (*fn)(), struct Iter out) +{ *out.point() = fn (); } + +extern "C" void __gxx_personality_v0 () +{ +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959.C b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C new file mode 100644 index 00000000000..169bf4aad17 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959.C @@ -0,0 +1,31 @@ +// { dg-additional-sources "pr71959-a.C" } + +// pr lto/71959 ICEd LTO due to mismatch between writing & reading behaviour + +struct Iter +{ + int *cursor; + + Iter(int *cursor_) : cursor(cursor_) {} + + int *point() const { return cursor; } +}; + +#pragma acc routine seq +int one () { return 1; } + +struct Apply +{ + static void apply (int (*fn)(), Iter out) + { *out.point() = fn (); } +}; + +int main () +{ + int x; + +#pragma acc parallel copyout(x) + Apply::apply (one, Iter (&x)); + + return x != 1; +} -- 2.17.1