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