Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper

[Ilya Verbin <iverbin@gmail.com>]

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 <iverbin@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