From: Cesar Philippidis <cesar@codesourcery.com>
To: "gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>
Subject: [patch,openacc] Fix PR71959: lto dump of callee counts
Date: Thu, 20 Sep 2018 18:06:00 -0000 [thread overview]
Message-ID: <319b3ebd-c601-449b-718c-963b68414224@codesourcery.com> (raw)
[-- Attachment #1: Type: text/plain, Size: 505 bytes --]
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
[-- Attachment #2: 0001-PR71959-lto-dump-of-callee-counts.patch --]
[-- Type: text/x-patch, Size: 3774 bytes --]
[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
next reply other threads:[~2018-09-20 17:55 UTC|newest]
Thread overview: 9+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-09-20 18:06 Cesar Philippidis [this message]
2018-09-25 13:01 ` Martin Jambor
2018-12-21 3:45 ` Julian Brown
2018-12-21 13:29 ` Julian Brown
2018-12-21 13:32 ` Jakub Jelinek
2018-12-21 17:09 ` Julian Brown
2018-12-22 18:53 ` Iain Sandoe
2019-01-08 13:31 ` Julian Brown
2019-01-08 13:36 ` Jakub Jelinek
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=319b3ebd-c601-449b-718c-963b68414224@codesourcery.com \
--to=cesar@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).