[openacc] Fix PR71959: lto dump of callee counts

Message ID 319b3ebd-c601-449b-718c-963b68414224@codesourcery.com
State New
Headers show
Series
  • [openacc] Fix PR71959: lto dump of callee counts
Related show

Commit Message

Cesar Philippidis Sept. 20, 2018, 5:54 p.m.
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

Comments

Martin Jambor Sept. 25, 2018, 12:59 p.m. | #1
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);
Julian Brown Dec. 21, 2018, 2:56 a.m. | #2
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
Julian Brown Dec. 21, 2018, 1:23 p.m. | #3
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;
+}
Jakub Jelinek Dec. 21, 2018, 1:31 p.m. | #4
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
Julian Brown Dec. 21, 2018, 4:47 p.m. | #5
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
Iain Sandoe Dec. 22, 2018, 3:09 p.m. | #6
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
Julian Brown Jan. 8, 2019, 1:30 p.m. | #7
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;
+}
Jakub Jelinek Jan. 8, 2019, 1:35 p.m. | #8
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

Patch

[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