On Sat, 22 Dec 2018 15:09:34 +0000 Iain Sandoe <idsan...@googlemail.com> wrote:
> Hi Julian, > > > On 21 Dec 2018, at 16:47, Julian Brown <jul...@codesourcery.com> > > wrote: > > > On Fri, 21 Dec 2018 14:31:19 +0100 > > Jakub Jelinek <ja...@redhat.com> wrote: > > > >> On Fri, Dec 21, 2018 at 01:23:03PM +0000, Julian Brown wrote: > >>> 2018-xx-yy Nathan Sidwell <nat...@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 <jul...@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; +}