On Fri, Feb 19, 2016 at 20:41:58 +0100, Thomas Schwinge wrote:
> Hi!
>
> On Thu, 2 Oct 2014 19:14:57 +0400, Ilya Verbin <iver...@gmail.com> wrote:
> > With this patch lto-wrapper performs invocation of mkoffload tool for each
> > offload target. This tool [...]
> > will compile IR from .gnu.offload_lto_* sections into offload
> > target code and embed the resultant code (offload image) into the new host's
> > object file.
>
> Consider the following scenario:
>
> $ cat < CSTS-214-acc.c
> int acc (void)
> {
> int a;
>
> #pragma acc parallel num_gangs (1) copyout (a)
> a = 100;
>
> return a;
> }
> $ cat < CSTS-214-test.c
> extern int acc (void);
>
> int main (void)
> {
> if (acc () != 100)
> __builtin_abort ();
>
> return 0;
> }
>
> Compile these two files as follows:
>
> $ [GCC] -fopenacc -c CSTS-214-acc.c
> $ x86_64-linux-gnu-ar -cr CSTS-214-acc.a CSTS-214-acc.o
> $ [GCC] -fopenacc CSTS-214-test.c CSTS-214-acc.a
>
> The last step will fail -- with incomprehensible diagnostics, ;-) as so
> often when offloading fails... Here's what's going on: the
> LTO/offloading machinery correctly identifies that it needs to process
> the CSTS-214-acc.c:acc function, present in the CSTS-214-acc.a archive
> file at a certain offset, and it "encodes" that as follows:
> CSTS-214-acc.a@0x9e (see lto-plugin/lto-plugin.c:claim_file_handler, the
> "file->offset != 0" code right at the beginning). This makes its way
> down through here:
>
> > --- a/gcc/lto-wrapper.c
> > +++ b/gcc/lto-wrapper.c
>
> > +/* Copy a file from SRC to DEST. */
> > +
> > +static void
> > +copy_file (const char *dest, const char *src)
> > +{
> > + [...]
> > +}
>
> > @@ -624,6 +852,54 @@ run_gcc (unsigned argc, char *argv[])
>
> > + /* If object files contain offload sections, but do not contain LTO
> > sections,
> > + then there is no need to perform a link-time recompilation, i.e.
> > + lto-wrapper is used only for a compilation of offload images. */
> > + if (have_offload && !have_lto)
> > + {
> > + for (i = 1; i < argc; ++i)
> > + if ([...])
> > + {
> > + char *out_file;
> > + /* Can be ".o" or ".so". */
> > + char *ext = strrchr (argv[i], '.');
> > + if (ext == NULL)
> > + out_file = make_temp_file ("");
> > + else
> > + out_file = make_temp_file (ext);
> > + /* The linker will delete the files we give it, so make copies. */
> > + copy_file (out_file, argv[i]);
> > + printf ("%s\n", out_file);
> > + }
> > +[...]
> > + goto finish;
> > + }
> > +
> > if (lto_mode == LTO_MODE_LTO)
> > {
> > flto_out = make_temp_file (".lto.o");
> > @@ -850,6 +1126,10 @@ cont:
> > obstack_free (&env_obstack, NULL);
> > }
> >
> > + finish:
> > + if (offloadend)
> > + printf ("%s\n", offloadend);
> > +
> > obstack_free (&argv_obstack, NULL);
> > }
>
> When we hit this, for argv "CSTS-214-acc.a@0x9e", the copy_file call will
> fail -- there is no "CSTS-214-acc.a@0x9e" file to copy. If we strip off
> the "@0x[...]" suffix (but still printf the filename including the
> suffix), then things work. I copied that bit of code from earlier in
> this function, where the same archive offset handling needs to be done.
> Probably that code should be refactored a bit.
>
> Also, I wonder if the "ext == NULL" case can really happen, and needs to
> be handled as done in the code cited above, or if that can be simplified?
> (Not yet tested that.)
>
> Will something like the following be OK to fix this issue, or is that
> something "that should not happen", should be fixed differently?
>
> --- gcc/lto-wrapper.c
> +++ gcc/lto-wrapper.c
> @@ -1161,15 +1161,31 @@ run_gcc (unsigned argc, char *argv[])
> && strncmp (argv[i], "-flinker-output=",
> sizeof ("-flinker-output=") - 1) != 0)
> {
> + char *p;
> + off_t file_offset = 0;
> + long loffset;
> + int consumed;
> + char *filename = argv[i];
> +
> + if ((p = strrchr (argv[i], '@'))
> + && p != argv[i]
> + && sscanf (p, "@%li%n", &loffset, &consumed) >= 1
> + && strlen (p) == (unsigned int) consumed)
> + {
> + filename = XNEWVEC (char, p - argv[i] + 1);
> + memcpy (filename, argv[i], p - argv[i]);
> + filename[p - argv[i]] = '\0';
> + file_offset = (off_t) loffset;
> + }
> +
> char *out_file;
> - /* Can be ".o" or ".so". */
> - char *ext = strrchr (argv[i], '.');
> + char *ext = strrchr (filename, '.');
> if (ext == NULL)
> out_file = make_temp_file ("");
> else
> out_file = make_temp_file (ext);
> /* The linker will delete the files we give it, so make copies. */
> - copy_file (out_file, argv[i]);
> + copy_file (out_file, filename);
> printf ("%s\n", out_file);
> }
> goto finish;
Yes, this part of lto-wrapper is awfully. This patch completely reworks it:
https://gcc.gnu.org/ml/gcc-patches/2016-02/msg00709.html
It's not yet fully ready, I'm going to send the final patch for review tomorrow.
-- Ilya
