On Mon, 18 Jan 2016, Ilya Verbin wrote:
> 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?)
Can lto-wrapper add them as plugin outputs? Or does that wreck ordering?
Richard.
>
> > > 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
>
>
--
Richard Biener <rguent...@suse.de>
SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB
21284 (AG Nuernberg)
