From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 87287 invoked by alias); 20 Sep 2018 17:55:06 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 86857 invoked by uid 89); 20 Sep 2018 17:55:05 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.5 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS autolearn=ham version=3.3.2 spammy=bodys, body's, Hx-languages-length:4162 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 20 Sep 2018 17:55:03 +0000 Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1g33AL-00053j-EZ from Cesar_Philippidis@mentor.com for gcc-patches@gcc.gnu.org; Thu, 20 Sep 2018 10:55:01 -0700 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Thu, 20 Sep 2018 10:54:59 -0700 From: Cesar Philippidis Subject: [patch,openacc] Fix PR71959: lto dump of callee counts To: "gcc-patches@gcc.gnu.org" Message-ID: <319b3ebd-c601-449b-718c-963b68414224@codesourcery.com> Date: Thu, 20 Sep 2018 18:06:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------D96F4543E134A0BE833B22A1" X-SW-Source: 2018-09/txt/msg01184.txt.bz2 --------------D96F4543E134A0BE833B22A1 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 505 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 --------------D96F4543E134A0BE833B22A1 Content-Type: text/x-patch; name="0001-PR71959-lto-dump-of-callee-counts.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0001-PR71959-lto-dump-of-callee-counts.patch" Content-length: 3774 [PR71959] lto dump of callee counts 2018-XX-YY Nathan Sidwell Cesar Philippidis 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 --------------D96F4543E134A0BE833B22A1--