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

[Thomas Schwinge <thomas@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 4714 bytes --]

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;


Grüße
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]