This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[patch,openacc] Fix PR71959: lto dump of callee counts
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>
- Date: Thu, 20 Sep 2018 10:54:58 -0700
- Subject: [patch,openacc] Fix PR71959: lto dump of callee counts
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.
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