On Fri, Jan 15, 2016 at 09:15:01 +0100, Richard Biener wrote:
> On Fri, 15 Jan 2016, Ilya Verbin wrote:
> > II) The __offload_func_table, __offload_funcs_end, __offload_var_table,
> > __offload_vars_end are now provided by the linker script, instead of
> > crtoffload{begin,end}.o, this allows to surround all offload objects, even
> > those that are not claimed by lto-plugin.
> > Unfortunately it works only with ld, but doen't work with gold, because
> > https://sourceware.org/bugzilla/show_bug.cgi?id=15373
> > Any thoughts how to enable this linker script for gold?
>
> The easiest way would probably to add this handling to the default
> "linker script" in gold. I don't see an easy way around requiring
> changes to gold here - maybe dumping the default linker script from
> bfd and injecting the rules with some scripting so you have a complete
> script. Though likely gold won't grok that result.
>
> Really a question for Ian though.
Or the gcc driver can add crtoffload{begin,end}.o, but the problem is that it
can't determine whether the program contains offloading or not. So it can add
them to all -fopenmp/-fopenacc programs, if the compiler was configured with
--enable-offload-targets=... The overhead would be about 340 bytes for
binaries which doesn't use offloading. Is this acceptable? (Jakub?)
> > I used the following testcase:
> > $ cat main.c
> > void foo1 ();
> > void foo2 ();
> > void foo3 ();
> > void foo4 ();
> >
> > int main ()
> > {
> > foo1 ();
> > foo2 ();
> > foo3 ();
> > foo4 ();
> > return 0;
> > }
> >
> > $ cat test.c
> > #include <stdio.h>
> > #include <omp.h>
> > #define MAKE_FN_NAME(x) foo ## x
> > #define FN_NAME(x) MAKE_FN_NAME(x)
> > void FN_NAME(NUM) ()
> > {
> > int x, d;
> > #pragma omp target map(from: x, d)
> > {
> > x = NUM;
> > d = omp_is_initial_device ();
> > }
> > printf ("%s:\t%s ()\tx = %d\n", d ? "HOST" : "TARGET", __FUNCTION__, x);
> > if (x != NUM)
> > printf ("--------^\n");
> > }
> >
> > $ gcc -DNUM=1 -c -flto test.c -o obj1.o
> > $ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
> > $ gcc -DNUM=3 -c test.c -o obj3.o
> > $ gcc -DNUM=4 -c -flto -fopenmp test.c -o obj4.o
> > $ gcc -c main.c -o main.o
> > $ gcc -fopenmp obj1.o obj2.o obj3.o obj4.o main.o && ./a.out
> > $ gcc -fopenmp obj2.o obj3.o obj4.o obj1.o main.o && ./a.out
> > $ gcc -fopenmp obj3.o obj1.o obj2.o obj4.o main.o && ./a.out
>
> Did you try linking an archive with both offload-but-no-LTO and
> offload-and-LTO objects inside?
No. And it didn't work, because archives are handled by ld a bit differently.
I will fix it. Thanks! From ld/ldlang.c:
/* Find the insert point for the plugin's replacement files. We
place them after the first claimed real object file, or if the
first claimed object is an archive member, after the last real
object file immediately preceding the archive.
-- Ilya
