From patchwork Thu Sep 20 17:54:58 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 972622 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-486081-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="eOl119DC"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 42GPX31YDXz9s8T for ; Fri, 21 Sep 2018 03:55:13 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:message-id:date:mime-version:content-type; q=dns; s= default; b=L1UKg0dCh32VwM9gsmc9sEKmhFTtSyFoPRxea8aJvOAhjtkPMks/4 rvwjKgN9u59Ps0SKz6Jbl7cZ8IeV/fCtsn6VKAe3wQ0fGw2CKzVktqhsGMr/bRuE c7fmUa3uEEk9rJF7ijGpZi5a5pW7YqDYtCGRlntPR/XqopP1Jd4x3E= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:message-id:date:mime-version:content-type; s= default; bh=wD9RIWlvzJdkxrS1P7r0ayXipok=; b=eOl119DCJPr+efzFosXW 0PR1QYYynE/l4/bynPMIRT50wn31I+jjNRfSrd1oLcPC5pSdU0FHGW+MOwX5Yt9m XqOqxkJ0rVCPTLkN88Q9f6WxDk5bjObJeJPkzUld6ue5oXj+1AN6+rrNdTKr0HeD N+KDQXIOOOyLxB0hZGAiaLc= 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-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@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 10:54:58 -0700 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 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 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