public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [patch,openacc] Fix PR71959: lto dump of callee counts
@ 2018-09-20 18:06 Cesar Philippidis
  2018-09-25 13:01 ` Martin Jambor
  0 siblings, 1 reply; 9+ messages in thread
From: Cesar Philippidis @ 2018-09-20 18:06 UTC (permalink / raw)
  To: gcc-patches

[-- 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


^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2018-09-20 18:06 [patch,openacc] Fix PR71959: lto dump of callee counts Cesar Philippidis
@ 2018-09-25 13:01 ` Martin Jambor
  2018-12-21  3:45   ` Julian Brown
  0 siblings, 1 reply; 9+ messages in thread
From: Martin Jambor @ 2018-09-25 13:01 UTC (permalink / raw)
  To: Cesar Philippidis, gcc-patches

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);

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2018-09-25 13:01 ` Martin Jambor
@ 2018-12-21  3:45   ` Julian Brown
  2018-12-21 13:29     ` Julian Brown
  0 siblings, 1 reply; 9+ messages in thread
From: Julian Brown @ 2018-12-21  3:45 UTC (permalink / raw)
  To: Martin Jambor; +Cc: Cesar Philippidis, gcc-patches

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

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2018-12-21  3:45   ` Julian Brown
@ 2018-12-21 13:29     ` Julian Brown
  2018-12-21 13:32       ` Jakub Jelinek
  0 siblings, 1 reply; 9+ messages in thread
From: Julian Brown @ 2018-12-21 13:29 UTC (permalink / raw)
  To: Martin Jambor; +Cc: Cesar Philippidis, gcc-patches, Thomas Schwinge

[-- Attachment #1: Type: text/plain, Size: 1437 bytes --]

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.

[-- Attachment #2: pr71959-test-only.diff --]
[-- Type: text/x-patch, Size: 1809 bytes --]

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;
+}

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2018-12-21 13:29     ` Julian Brown
@ 2018-12-21 13:32       ` Jakub Jelinek
  2018-12-21 17:09         ` Julian Brown
  0 siblings, 1 reply; 9+ messages in thread
From: Jakub Jelinek @ 2018-12-21 13:32 UTC (permalink / raw)
  To: Julian Brown
  Cc: Martin Jambor, Cesar Philippidis, gcc-patches, Thomas Schwinge

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

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2018-12-21 13:32       ` Jakub Jelinek
@ 2018-12-21 17:09         ` Julian Brown
  2018-12-22 18:53           ` Iain Sandoe
  0 siblings, 1 reply; 9+ messages in thread
From: Julian Brown @ 2018-12-21 17:09 UTC (permalink / raw)
  To: Jakub Jelinek
  Cc: Martin Jambor, Cesar Philippidis, gcc-patches, Thomas Schwinge

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

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2018-12-21 17:09         ` Julian Brown
@ 2018-12-22 18:53           ` Iain Sandoe
  2019-01-08 13:31             ` Julian Brown
  0 siblings, 1 reply; 9+ messages in thread
From: Iain Sandoe @ 2018-12-22 18:53 UTC (permalink / raw)
  To: Julian Brown
  Cc: Jakub Jelinek, Martin Jambor, Cesar Philippidis, gcc-patches,
	Thomas Schwinge

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

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2018-12-22 18:53           ` Iain Sandoe
@ 2019-01-08 13:31             ` Julian Brown
  2019-01-08 13:36               ` Jakub Jelinek
  0 siblings, 1 reply; 9+ messages in thread
From: Julian Brown @ 2019-01-08 13:31 UTC (permalink / raw)
  To: Iain Sandoe
  Cc: Jakub Jelinek, Martin Jambor, Cesar Philippidis, gcc-patches,
	Thomas Schwinge

[-- Attachment #1: Type: text/plain, Size: 1847 bytes --]

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

[-- Attachment #2: pr71959-test-only-3.diff --]
[-- Type: text/x-patch, Size: 1967 bytes --]

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;
+}

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [patch,openacc] Fix PR71959: lto dump of callee counts
  2019-01-08 13:31             ` Julian Brown
@ 2019-01-08 13:36               ` Jakub Jelinek
  0 siblings, 0 replies; 9+ messages in thread
From: Jakub Jelinek @ 2019-01-08 13:36 UTC (permalink / raw)
  To: Julian Brown
  Cc: Iain Sandoe, Martin Jambor, Cesar Philippidis, gcc-patches,
	Thomas Schwinge

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

^ permalink raw reply	[flat|nested] 9+ messages in thread

end of thread, other threads:[~2019-01-08 13:36 UTC | newest]

Thread overview: 9+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2018-09-20 18:06 [patch,openacc] Fix PR71959: lto dump of callee counts Cesar Philippidis
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

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).