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

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

[Ilya Verbin]

Hello,

With this patch lto-wrapper performs invocation of mkoffload tool for each
offload target.  This tool should be provided by the corresponding offload
compiler.  It 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.

Bootstrapped and regtested on top of patch 3.  Is it OK for trunk?

Thanks,
  -- Ilya


2014-10-02  Ilya Verbin  <ilya.verbin@intel.com>
	    Bernd Schmidt  <bernds@codesourcery.com>
	    Andrey Turetskiy  <andrey.turetskiy@intel.com>
	    Michael Zolotukhin  <michael.v.zolotukhin@intel.com>

gcc/
	* gcc.c (spec_host_machine, accel_dir_suffix): New variables.
	(process_command): Tweak path construction for the possibility
	of being configured as an offload compiler.
	(main): Likewise.  Look up specs in just_machine_suffix only if not
	ACCEL_COMPILER.  Construct OFFLOAD_TARGET_NAMES environment variable if
	we have OFFLOAD_TARGETS.
	* lto-wrapper.c (OFFLOAD_TARGET_NAMES_ENV): Define.
	(offload_names, offloadbegin, offloadend): New static variables.
	(free_array_of_ptrs, parse_env_var, access_check, compile_offload_image)
	(compile_images_for_offload_targets, copy_file, find_offloadbeginend):
	New static functions.
	(run_gcc): Determine whether offload sections are present.  If so, run
	compile_images_for_offload_targets and return the names of new generated
	objects to linker.  If there are offload sections, but no LTO sections,
	then return the copies of input objects without link-time recompilation.
lto-plugin/
	* lto-plugin.c (OFFLOAD_SECTION, OFFLOAD_SECTION_LEN): Define.
	(struct plugin_objfile): Add new field "offload".
	(process_offload_section): New static function.
	(claim_file_handler): Claim file if it contains offload sections.

---

diff --git a/gcc/gcc.c b/gcc/gcc.c
index 47c4e28..82509a8 100644
--- a/gcc/gcc.c
+++ b/gcc/gcc.c
@@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION;
 /* The target machine.  */
 
 static const char *spec_machine = DEFAULT_TARGET_MACHINE;
+static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE;
 
 /* Nonzero if cross-compiling.
    When -b is used, the value comes from the `specs' file.  */
@@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2
    relative to the driver.  */
 static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX;
 
+/* A prefix to be used when this is an accelerator compiler.  */
+static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX;
+
 /* Subdirectory to use for locating libraries.  Set by
    set_multilib_dir based on the compilation options.  */
 
@@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count,
     }
 
   gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
-  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
+  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
 			    dir_separator_str, NULL);
 
   /* Look for tools relative to the location from which the driver is
      running, or, if that is not available, the configured prefix.  */
   tooldir_prefix
     = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix,
-	      spec_machine, dir_separator_str,
-	      spec_version, dir_separator_str, tooldir_prefix2, NULL);
+	      spec_host_machine, dir_separator_str, spec_version,
+	      accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL);
   free (tooldir_prefix2);
 
   add_prefix (&exec_prefixes,
@@ -6878,8 +6882,8 @@ main (int argc, char **argv)
 
   /* Read specs from a file if there is one.  */
 
-  machine_suffix = concat (spec_machine, dir_separator_str,
-			   spec_version, dir_separator_str, NULL);
+  machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version,
+			   accel_dir_suffix, dir_separator_str, NULL);
   just_machine_suffix = concat (spec_machine, dir_separator_str, NULL);
 
   specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true);
@@ -6889,16 +6893,17 @@ main (int argc, char **argv)
   else
     init_spec ();
 
+#ifndef ACCEL_COMPILER
   /* We need to check standard_exec_prefix/just_machine_suffix/specs
      for any override of as, ld and libraries.  */
   specs_file = (char *) alloca (strlen (standard_exec_prefix)
 		       + strlen (just_machine_suffix) + sizeof ("specs"));
-
   strcpy (specs_file, standard_exec_prefix);
   strcat (specs_file, just_machine_suffix);
   strcat (specs_file, "specs");
   if (access (specs_file, R_OK) == 0)
     read_specs (specs_file, true, false);
+#endif
 
   /* Process any configure-time defaults specified for the command line
      options, via OPTION_DEFAULT_SPECS.  */
@@ -7077,8 +7082,9 @@ main (int argc, char **argv)
 
   /* If we have a GCC_EXEC_PREFIX envvar, modify it for cpp's sake.  */
   if (gcc_exec_prefix)
-    gcc_exec_prefix = concat (gcc_exec_prefix, spec_machine, dir_separator_str,
-			      spec_version, dir_separator_str, NULL);
+    gcc_exec_prefix = concat (gcc_exec_prefix, spec_host_machine,
+			      dir_separator_str, spec_version,
+			      accel_dir_suffix, dir_separator_str, NULL);
 
   /* Now we have the specs.
      Set the `valid' bits for switches that match anything in any spec.  */
@@ -7097,6 +7103,16 @@ main (int argc, char **argv)
   obstack_grow (&collect_obstack, argv[0], strlen (argv[0]) + 1);
   xputenv (XOBFINISH (&collect_obstack, char *));
 
+  if (strlen (OFFLOAD_TARGETS) > 0)
+    {
+      obstack_init (&collect_obstack);
+      obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=",
+		    sizeof ("OFFLOAD_TARGET_NAMES=") - 1);
+      obstack_grow (&collect_obstack, OFFLOAD_TARGETS,
+		    strlen (OFFLOAD_TARGETS) + 1);
+      xputenv (XOBFINISH (&collect_obstack, char *));
+    }
+
   /* Set up to remember the pathname of the lto wrapper. */
 
   if (have_c)
diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c
index 08fd090..2c9b503 100644
--- a/gcc/lto-wrapper.c
+++ b/gcc/lto-wrapper.c
@@ -49,6 +49,8 @@ along with GCC; see the file COPYING3.  If not see
 #include "lto-section-names.h"
 #include "collect-utils.h"
 
+#define OFFLOAD_TARGET_NAMES_ENV	"OFFLOAD_TARGET_NAMES"
+
 enum lto_mode_d {
   LTO_MODE_NONE,			/* Not doing LTO.  */
   LTO_MODE_LTO,				/* Normal LTO.  */
@@ -63,6 +65,8 @@ static char *flto_out;
 static unsigned int nr;
 static char **input_names;
 static char **output_names;
+static char **offload_names;
+static const char *offloadbegin, *offloadend;
 static char *makefile;
 
 const char tool_name[] = "lto-wrapper";
@@ -358,6 +362,224 @@ merge_and_complain (struct cl_decoded_option **decoded_options,
     }
 }
 
+/* Auxiliary function that frees elements of PTR and PTR itself.
+   N is number of elements to be freed.  If PTR is NULL, nothing is freed.
+   If an element is NULL, subsequent elements are not freed.  */
+
+static void **
+free_array_of_ptrs (void **ptr, unsigned n)
+{
+  if (!ptr)
+    return NULL;
+  for (unsigned i = 0; i < n; i++)
+    {
+      if (!ptr[i])
+	break;
+      free (ptr[i]);
+    }
+  free (ptr);
+  return NULL;
+}
+
+/* Parse STR, saving found tokens into PVALUES and return their number.
+   Tokens are assumed to be delimited by ':'.  If APPEND is non-null,
+   append it to every token we find.  */
+
+static unsigned
+parse_env_var (const char *str, char ***pvalues, const char *append)
+{
+  const char *curval, *nextval;
+  char **values;
+  unsigned num = 1, i;
+
+  curval = strchr (str, ':');
+  while (curval)
+    {
+      num++;
+      curval = strchr (curval + 1, ':');
+    }
+
+  values = (char**) xmalloc (num * sizeof (char*));
+  curval = str;
+  nextval = strchrnul (curval, ':');
+
+  int append_len = append ? strlen (append) : 0;
+  for (i = 0; i < num; i++)
+    {
+      int l = nextval - curval;
+      values[i] = (char*) xmalloc (l + 1 + append_len);
+      memcpy (values[i], curval, l);
+      values[i][l] = 0;
+      if (append)
+	strcat (values[i], append);
+      curval = nextval + 1;
+      nextval = strchrnul (curval, ':');
+    }
+  *pvalues = values;
+  return num;
+}
+
+/* Check whether NAME can be accessed in MODE.  This is like access,
+   except that it never considers directories to be executable.  */
+
+static int
+access_check (const char *name, int mode)
+{
+  if (mode == X_OK)
+    {
+      struct stat st;
+
+      if (stat (name, &st) < 0
+	  || S_ISDIR (st.st_mode))
+	return -1;
+    }
+
+  return access (name, mode);
+}
+
+/* Prepare a target image for offload TARGET, using mkoffload tool from
+   COMPILER_PATH.  Return the name of the resultant object file.  */
+
+static char *
+compile_offload_image (const char *target, const char *compiler_path,
+		       unsigned in_argc, char *in_argv[])
+{
+  char *filename = NULL;
+  char **argv;
+  char *suffix
+    = XALLOCAVEC (char, strlen ("/accel//mkoffload") + strlen (target) + 1);
+  strcpy (suffix, "/accel/");
+  strcat (suffix, target);
+  strcat (suffix, "/mkoffload");
+
+  char **paths = NULL;
+  unsigned n_paths = parse_env_var (compiler_path, &paths, suffix);
+
+  const char *compiler = NULL;
+  for (unsigned i = 0; i < n_paths; i++)
+    if (access_check (paths[i], X_OK) == 0)
+      {
+	compiler = paths[i];
+	break;
+      }
+
+  if (!compiler)
+    goto out;
+
+  /* Generate temporary output file name.  */
+  filename = make_temp_file (".target.o");
+
+  struct obstack argv_obstack;
+  obstack_init (&argv_obstack);
+  obstack_ptr_grow (&argv_obstack, compiler);
+  obstack_ptr_grow (&argv_obstack, "-o");
+  obstack_ptr_grow (&argv_obstack, filename);
+
+  for (unsigned i = 1; i < in_argc; i++)
+    obstack_ptr_grow (&argv_obstack, in_argv[i]);
+  obstack_ptr_grow (&argv_obstack, NULL);
+
+  argv = XOBFINISH (&argv_obstack, char **);
+  fork_execute (argv[0], argv, true);
+  obstack_free (&argv_obstack, NULL);
+
+ out:
+  free_array_of_ptrs ((void **) paths, n_paths);
+  return filename;
+}
+
+
+/* The main routine dealing with offloading.
+   The routine builds a target image for each offload target.  IN_ARGC and
+   IN_ARGV specify options and input object files.  As all of them could contain
+   target sections, we pass them all to target compilers.  */
+
+static void
+compile_images_for_offload_targets (unsigned in_argc, char *in_argv[])
+{
+  char **names = NULL;
+  const char *target_names = getenv (OFFLOAD_TARGET_NAMES_ENV);
+  if (!target_names)
+    return;
+  unsigned num_targets = parse_env_var (target_names, &names, NULL);
+
+  const char *compiler_path = getenv ("COMPILER_PATH");
+  if (!compiler_path)
+    goto out;
+
+  /* Prepare an image for each target and save the name of the resultant object
+     file to the OFFLOAD_NAMES array.  It is terminated by a NULL entry.  */
+  offload_names = XCNEWVEC (char *, num_targets + 1);
+  for (unsigned i = 0; i < num_targets; i++)
+    {
+      offload_names[i] = compile_offload_image (names[i], compiler_path,
+						in_argc, in_argv);
+      if (!offload_names[i])
+	fatal_error ("problem with building target image for %s\n", names[i]);
+    }
+
+ out:
+  free_array_of_ptrs ((void **) names, num_targets);
+}
+
+/* Copy a file from SRC to DEST.  */
+
+static void
+copy_file (const char *dest, const char *src)
+{
+  FILE *d = fopen (dest, "wb");
+  FILE *s = fopen (src, "rb");
+  char buffer[512];
+  while (!feof (s))
+    {
+      size_t len = fread (buffer, 1, 512, s);
+      if (ferror (s) != 0)
+	fatal_error ("reading input file");
+      if (len > 0)
+	{
+	  fwrite (buffer, 1, len, d);
+	  if (ferror (d) != 0)
+	    fatal_error ("writing output file");
+	}
+    }
+}
+
+/* Find the crtoffloadbegin.o and crtoffloadend.o files in LIBRARY_PATH, make
+   copies and store the names of the copies in offloadbegin and offloadend.  */
+
+static void
+find_offloadbeginend (void)
+{
+  char **paths = NULL;
+  const char *library_path = getenv ("LIBRARY_PATH");
+  if (!library_path)
+    return;
+  unsigned n_paths = parse_env_var (library_path, &paths, "/crtoffloadbegin.o");
+
+  unsigned i;
+  for (i = 0; i < n_paths; i++)
+    if (access_check (paths[i], R_OK) == 0)
+      {
+	size_t len = strlen (paths[i]);
+	char *tmp = xstrdup (paths[i]);
+	strcpy (paths[i] + len - strlen ("begin.o"), "end.o");
+	if (access_check (paths[i], R_OK) != 0)
+	  fatal_error ("installation error, can't find crtoffloadend.o");
+	/* The linker will delete the filenames we give it, so make
+	   copies.  */
+	offloadbegin = make_temp_file (".o");
+	offloadend = make_temp_file (".o");
+	copy_file (offloadbegin, tmp);
+	copy_file (offloadend, paths[i]);
+	free (tmp);
+	break;
+      }
+  if (i == n_paths)
+    fatal_error ("installation error, can't find crtoffloadbegin.o");
+
+  free_array_of_ptrs ((void **) paths, n_paths);
+}
+
 /* Execute gcc. ARGC is the number of arguments. ARGV contains the arguments. */
 
 static void
@@ -378,6 +600,8 @@ run_gcc (unsigned argc, char *argv[])
   unsigned int decoded_options_count;
   struct obstack argv_obstack;
   int new_head_argc;
+  bool have_lto = false;
+  bool have_offload = false;
 
   /* Get the driver and options.  */
   collect_gcc = getenv ("COLLECT_GCC");
@@ -426,6 +650,9 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
+      if (simple_object_find_section (sobj, OFFLOAD_SECTION_NAME_PREFIX ".opts",
+				      &offset, &length, &errmsg, &err))
+	have_offload = true;
       if (!simple_object_find_section (sobj, LTO_SECTION_NAME_PREFIX "." "opts",
 				       &offset, &length, &errmsg, &err))
 	{
@@ -433,6 +660,7 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
+      have_lto = true;
       lseek (fd, file_offset + offset, SEEK_SET);
       data = (char *)xmalloc (length);
       read (fd, data, length);
@@ -624,6 +852,54 @@ run_gcc (unsigned argc, char *argv[])
   /* Remember at which point we can scrub args to re-use the commons.  */
   new_head_argc = obstack_object_size (&argv_obstack) / sizeof (void *);
 
+  if (have_offload)
+    {
+      compile_images_for_offload_targets (argc, argv);
+      if (offload_names)
+	{
+	  find_offloadbeginend ();
+	  for (i = 0; offload_names[i]; i++)
+	    printf ("%s\n", offload_names[i]);
+	  free_array_of_ptrs ((void **) offload_names, i);
+	}
+    }
+
+  if (offloadbegin)
+    printf ("%s\n", offloadbegin);
+
+  /* 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 (strncmp (argv[i], "-fresolution=", sizeof ("-fresolution=") - 1))
+	  {
+	    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);
+	  }
+
+      /* By default linker does not discard .gnu.offload_lto_* sections.  */
+      const char *linker_script = make_temp_file ("_linker_script.x");
+      FILE *stream = fopen (linker_script, "w");
+      if (!stream)
+	fatal_error ("fopen %s: %m", linker_script);
+      fprintf (stream, "SECTIONS { /DISCARD/ : { *("
+		       OFFLOAD_SECTION_NAME_PREFIX "*) } }\n");
+      fclose (stream);
+      printf ("%s\n", linker_script);
+
+      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);
 }
 
diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c
index 910e23c..fb6555d 100644
--- a/lto-plugin/lto-plugin.c
+++ b/lto-plugin/lto-plugin.c
@@ -86,6 +86,8 @@ along with this program; see the file COPYING3.  If not see
 
 #define LTO_SECTION_PREFIX	".gnu.lto_.symtab"
 #define LTO_SECTION_PREFIX_LEN	(sizeof (LTO_SECTION_PREFIX) - 1)
+#define OFFLOAD_SECTION		".gnu.offload_lto_.opts"
+#define OFFLOAD_SECTION_LEN	(sizeof (OFFLOAD_SECTION) - 1)
 
 /* The part of the symbol table the plugin has to keep track of. Note that we
    must keep SYMS until all_symbols_read is called to give the linker time to
@@ -111,6 +113,7 @@ struct plugin_symtab
 struct plugin_objfile
 {
   int found;
+  int offload;
   simple_object_read *objfile;
   struct plugin_symtab *out;
   const struct ld_plugin_input_file *file;
@@ -862,6 +865,21 @@ err:
   return 0;
 }
 
+/* Find an offload section of an object file.  */
+
+static int
+process_offload_section (void *data, const char *name, off_t offset, off_t len)
+{
+  if (!strncmp (name, OFFLOAD_SECTION, OFFLOAD_SECTION_LEN))
+    {
+      struct plugin_objfile *obj = (struct plugin_objfile *) data;
+      obj->offload = 1;
+      return 0;
+    }
+
+  return 1;
+}
+
 /* Callback used by gold to check if the plugin will claim FILE. Writes
    the result in CLAIMED. */
 
@@ -899,6 +917,7 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed)
   *claimed = 0;
   obj.file = file;
   obj.found = 0;
+  obj.offload = 0;
   obj.out = &lto_file.symtab;
   errmsg = NULL;
   obj.objfile = simple_object_start_read (file->fd, file->offset, LTO_SEGMENT_NAME,
@@ -920,7 +939,11 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed)
       goto err;
     }
 
-  if (obj.found == 0)
+  if (obj.objfile)
+    simple_object_find_sections (obj.objfile, process_offload_section,
+				 &obj, &err);
+
+  if (obj.found == 0 && obj.offload == 0)
     goto err;
 
   if (obj.found > 1)
-- 
1.7.1

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

[Jakub Jelinek]

On Thu, Oct 02, 2014 at 07:14:57PM +0400, Ilya Verbin wrote:
> @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2
>     relative to the driver.  */
>  static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX;
>  
> +/* A prefix to be used when this is an accelerator compiler.  */
> +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX;

Is ACCEL_DIR_SUFFIX here "" or something starting with "/ ?

> @@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count,
>      }
>  
>    gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
> -  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
> +  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
>  			    dir_separator_str, NULL);
>  
>    /* Look for tools relative to the location from which the driver is
>       running, or, if that is not available, the configured prefix.  */
>    tooldir_prefix
>      = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix,
> -	      spec_machine, dir_separator_str,
> -	      spec_version, dir_separator_str, tooldir_prefix2, NULL);
> +	      spec_host_machine, dir_separator_str, spec_version,
> +	      accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL);
>    free (tooldir_prefix2);
>  
>    add_prefix (&exec_prefixes,

The reason I'm asking is that otherwise it seems gcc.c heavily uses
dir_separator_str for the separators.  Don't have any experience with
targets where DIR_SEPARATOR is not / though, perhaps it is a non-issue.

> @@ -6878,8 +6882,8 @@ main (int argc, char **argv)
>  
>    /* Read specs from a file if there is one.  */
>  
> -  machine_suffix = concat (spec_machine, dir_separator_str,
> -			   spec_version, dir_separator_str, NULL);
> +  machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version,
> +			   accel_dir_suffix, dir_separator_str, NULL);
>    just_machine_suffix = concat (spec_machine, dir_separator_str, NULL);
>  
>    specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true);
> @@ -6889,16 +6893,17 @@ main (int argc, char **argv)
>    else
>      init_spec ();
>  
> +#ifndef ACCEL_COMPILER
>    /* We need to check standard_exec_prefix/just_machine_suffix/specs
>       for any override of as, ld and libraries.  */
>    specs_file = (char *) alloca (strlen (standard_exec_prefix)
>  		       + strlen (just_machine_suffix) + sizeof ("specs"));
> -
>    strcpy (specs_file, standard_exec_prefix);
>    strcat (specs_file, just_machine_suffix);
>    strcat (specs_file, "specs");
>    if (access (specs_file, R_OK) == 0)
>      read_specs (specs_file, true, false);
> +#endif

Why do you want to disable specs reading for the accel compiler?
Then users won't have the possibility to override defaults etc. easily...

> @@ -7097,6 +7103,16 @@ main (int argc, char **argv)
>    obstack_grow (&collect_obstack, argv[0], strlen (argv[0]) + 1);
>    xputenv (XOBFINISH (&collect_obstack, char *));
>  
> +  if (strlen (OFFLOAD_TARGETS) > 0)
> +    {
> +      obstack_init (&collect_obstack);
> +      obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=",
> +		    sizeof ("OFFLOAD_TARGET_NAMES=") - 1);
> +      obstack_grow (&collect_obstack, OFFLOAD_TARGETS,
> +		    strlen (OFFLOAD_TARGETS) + 1);
> +      xputenv (XOBFINISH (&collect_obstack, char *));
> +    }
> +

I'm surprised to see the obstack_init call here, but looking at
gcc.c, I'm surprised to see it in more places.

I've always thought that obstack_init was something you invoke generally
once on a given obstack object, then work with the obstack and then
obstack_free (..., NULL) it at the end.  Now, if it wants the obstack to be
live until exit, it just would not obstack_free it.  But calling
obstack_init on the already initialized obstack results IMHO in memory
leaks.  It should be initialized just once somewhere.

>    /* Set up to remember the pathname of the lto wrapper. */
>  
>    if (have_c)
> diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c
> index 08fd090..2c9b503 100644
> --- a/gcc/lto-wrapper.c
> +++ b/gcc/lto-wrapper.c
> @@ -49,6 +49,8 @@ along with GCC; see the file COPYING3.  If not see
>  #include "lto-section-names.h"
>  #include "collect-utils.h"
>  
> +#define OFFLOAD_TARGET_NAMES_ENV	"OFFLOAD_TARGET_NAMES"

Missing comment about the env var and what it is for.

> +/* Prepare a target image for offload TARGET, using mkoffload tool from
> +   COMPILER_PATH.  Return the name of the resultant object file.  */
> +
> +static char *
> +compile_offload_image (const char *target, const char *compiler_path,
> +		       unsigned in_argc, char *in_argv[])
> +{
> +  char *filename = NULL;
> +  char **argv;
> +  char *suffix
> +    = XALLOCAVEC (char, strlen ("/accel//mkoffload") + strlen (target) + 1);

Use sizeof ("/accel//mkoffload") + strlen (target) instead?

> +  strcpy (suffix, "/accel/");
> +  strcat (suffix, target);
> +  strcat (suffix, "/mkoffload");
> +
> +  char **paths = NULL;
> +  unsigned n_paths = parse_env_var (compiler_path, &paths, suffix);
> +
> +  const char *compiler = NULL;
> +  for (unsigned i = 0; i < n_paths; i++)
> +    if (access_check (paths[i], X_OK) == 0)
> +      {
> +	compiler = paths[i];
> +	break;
> +      }
> +
> +  if (!compiler)
> +    goto out;
> +
> +  /* Generate temporary output file name.  */
> +  filename = make_temp_file (".target.o");
> +
> +  struct obstack argv_obstack;
> +  obstack_init (&argv_obstack);
> +  obstack_ptr_grow (&argv_obstack, compiler);
> +  obstack_ptr_grow (&argv_obstack, "-o");
> +  obstack_ptr_grow (&argv_obstack, filename);
> +
> +  for (unsigned i = 1; i < in_argc; i++)
> +    obstack_ptr_grow (&argv_obstack, in_argv[i]);
> +  obstack_ptr_grow (&argv_obstack, NULL);
> +
> +  argv = XOBFINISH (&argv_obstack, char **);
> +  fork_execute (argv[0], argv, true);
> +  obstack_free (&argv_obstack, NULL);
> +
> + out:

The goto is probably unnecessary here, indenting all the lines by
4 more spaces and using if (compiler) { ... } instead doesn't need
any further warpping.

> +  free_array_of_ptrs ((void **) paths, n_paths);
> +  return filename;
> +}
> +

> +      /* By default linker does not discard .gnu.offload_lto_* sections.  */
> +      const char *linker_script = make_temp_file ("_linker_script.x");
> +      FILE *stream = fopen (linker_script, "w");
> +      if (!stream)
> +	fatal_error ("fopen %s: %m", linker_script);
> +      fprintf (stream, "SECTIONS { /DISCARD/ : { *("
> +		       OFFLOAD_SECTION_NAME_PREFIX "*) } }\n");
> +      fclose (stream);
> +      printf ("%s\n", linker_script);
> +
> +      goto finish;
> +    }

Does this work with gold?  Are there any other linkers that support plugins,
but don't support linker scripts this way?


	Jakub

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

[Ilya Verbin]

On 08 Oct 12:26, Jakub Jelinek wrote:
> On Thu, Oct 02, 2014 at 07:14:57PM +0400, Ilya Verbin wrote:
> > @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2
> >     relative to the driver.  */
> >  static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX;
> >  
> > +/* A prefix to be used when this is an accelerator compiler.  */
> > +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX;
> 
> Is ACCEL_DIR_SUFFIX here "" or something starting with "/ ?
> 
> The reason I'm asking is that otherwise it seems gcc.c heavily uses
> dir_separator_str for the separators.  Don't have any experience with
> targets where DIR_SEPARATOR is not / though, perhaps it is a non-issue.

It is "" for the host compiler and starts with "/" for the accel compiler.
Looking at gcc/configure.ac, there are a lot of paths with slashes.
So, I think it's ok to define ACCEL_DIR_SUFFIX in such a way.

> > +#ifndef ACCEL_COMPILER
> >    /* We need to check standard_exec_prefix/just_machine_suffix/specs
> >       for any override of as, ld and libraries.  */
> >    specs_file = (char *) alloca (strlen (standard_exec_prefix)
> >  		       + strlen (just_machine_suffix) + sizeof ("specs"));
> > -
> >    strcpy (specs_file, standard_exec_prefix);
> >    strcat (specs_file, just_machine_suffix);
> >    strcat (specs_file, "specs");
> >    if (access (specs_file, R_OK) == 0)
> >      read_specs (specs_file, true, false);
> > +#endif
> 
> Why do you want to disable specs reading for the accel compiler?
> Then users won't have the possibility to override defaults etc. easily...

Bernd,
Do you need this ifndef for PTX?  Or I could remove it?

> > @@ -7097,6 +7103,16 @@ main (int argc, char **argv)
> >    obstack_grow (&collect_obstack, argv[0], strlen (argv[0]) + 1);
> >    xputenv (XOBFINISH (&collect_obstack, char *));
> >  
> > +  if (strlen (OFFLOAD_TARGETS) > 0)
> > +    {
> > +      obstack_init (&collect_obstack);
> > +      obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=",
> > +		    sizeof ("OFFLOAD_TARGET_NAMES=") - 1);
> > +      obstack_grow (&collect_obstack, OFFLOAD_TARGETS,
> > +		    strlen (OFFLOAD_TARGETS) + 1);
> > +      xputenv (XOBFINISH (&collect_obstack, char *));
> > +    }
> > +
> 
> I'm surprised to see the obstack_init call here, but looking at
> gcc.c, I'm surprised to see it in more places.
> 
> I've always thought that obstack_init was something you invoke generally
> once on a given obstack object, then work with the obstack and then
> obstack_free (..., NULL) it at the end.  Now, if it wants the obstack to be
> live until exit, it just would not obstack_free it.  But calling
> obstack_init on the already initialized obstack results IMHO in memory
> leaks.  It should be initialized just once somewhere.

Right, I removed obstack_init from this patch.

> > +#define OFFLOAD_TARGET_NAMES_ENV	"OFFLOAD_TARGET_NAMES"
> 
> Missing comment about the env var and what it is for.

Fixed.

> > +  char *suffix
> > +    = XALLOCAVEC (char, strlen ("/accel//mkoffload") + strlen (target) + 1);
> 
> Use sizeof ("/accel//mkoffload") + strlen (target) instead?

Done.

> > + out:
> 
> The goto is probably unnecessary here, indenting all the lines by
> 4 more spaces and using if (compiler) { ... } instead doesn't need
> any further warpping.

Done.

> > +      /* By default linker does not discard .gnu.offload_lto_* sections.  */
> > +      const char *linker_script = make_temp_file ("_linker_script.x");
> > +      FILE *stream = fopen (linker_script, "w");
> > +      if (!stream)
> > +	fatal_error ("fopen %s: %m", linker_script);
> > +      fprintf (stream, "SECTIONS { /DISCARD/ : { *("
> > +		       OFFLOAD_SECTION_NAME_PREFIX "*) } }\n");
> > +      fclose (stream);
> > +      printf ("%s\n", linker_script);
> > +
> > +      goto finish;
> > +    }
> 
> Does this work with gold?  Are there any other linkers that support plugins,
> but don't support linker scripts this way?

Oops, gold does not support scripts, outputted from plugins :(
"error: SECTIONS seen after other input files; try -T/--script"

Probably, we should update default linker scripts in binutils?
But without latest ld/gold all binaries compiled without -flto and with
offloading will contain intermediate bytecode...

Thanks,
  -- Ilya

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

[Bernd Schmidt]

On 10/09/2014 02:07 PM, Ilya Verbin wrote:
>>> +#ifndef ACCEL_COMPILER
>>>     /* We need to check standard_exec_prefix/just_machine_suffix/specs
>>>        for any override of as, ld and libraries.  */
>>>     specs_file = (char *) alloca (strlen (standard_exec_prefix)
>>>   		       + strlen (just_machine_suffix) + sizeof ("specs"));
>>> -
>>>     strcpy (specs_file, standard_exec_prefix);
>>>     strcat (specs_file, just_machine_suffix);
>>>     strcat (specs_file, "specs");
>>>     if (access (specs_file, R_OK) == 0)
>>>       read_specs (specs_file, true, false);
>>> +#endif
>>
>> Why do you want to disable specs reading for the accel compiler?
>> Then users won't have the possibility to override defaults etc. easily...
>
> Bernd,
> Do you need this ifndef for PTX?  Or I could remove it?

I suspect the paths aren't right. If that can be fixed it should be fine 
to remove the ifdef.


Bernd

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

[Ilya Verbin]

On 09 Oct 16:07, Ilya Verbin wrote:
> > > +      /* By default linker does not discard .gnu.offload_lto_* sections.  */
> > > +      const char *linker_script = make_temp_file ("_linker_script.x");
> > > +      FILE *stream = fopen (linker_script, "w");
> > > +      if (!stream)
> > > +	fatal_error ("fopen %s: %m", linker_script);
> > > +      fprintf (stream, "SECTIONS { /DISCARD/ : { *("
> > > +		       OFFLOAD_SECTION_NAME_PREFIX "*) } }\n");
> > > +      fclose (stream);
> > > +      printf ("%s\n", linker_script);
> > > +
> > > +      goto finish;
> > > +    }
> > 
> > Does this work with gold?  Are there any other linkers that support plugins,
> > but don't support linker scripts this way?
> 
> Oops, gold does not support scripts, outputted from plugins :(
> "error: SECTIONS seen after other input files; try -T/--script"
> 
> Probably, we should update default linker scripts in binutils?
> But without latest ld/gold all binaries compiled without -flto and with
> offloading will contain intermediate bytecode...

Actually, this issue is not due to outputting a script from a plugin,
gold just does not support partial linker scripts:
https://sourceware.org/bugzilla/show_bug.cgi?id=17451

So it seems that discarding .gnu.offload_lto_* sections (like it is done for
.gnu.lto_*) in the default ld and gold scripts is the right way?

Thanks,
  -- Ilya

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

[Jakub Jelinek]

On Fri, Oct 10, 2014 at 12:07:03AM +0400, Ilya Verbin wrote:
> On 09 Oct 16:07, Ilya Verbin wrote:
> > > > +      /* By default linker does not discard .gnu.offload_lto_* sections.  */
> > > > +      const char *linker_script = make_temp_file ("_linker_script.x");
> > > > +      FILE *stream = fopen (linker_script, "w");
> > > > +      if (!stream)
> > > > +	fatal_error ("fopen %s: %m", linker_script);
> > > > +      fprintf (stream, "SECTIONS { /DISCARD/ : { *("
> > > > +		       OFFLOAD_SECTION_NAME_PREFIX "*) } }\n");
> > > > +      fclose (stream);
> > > > +      printf ("%s\n", linker_script);
> > > > +
> > > > +      goto finish;
> > > > +    }
> > > 
> > > Does this work with gold?  Are there any other linkers that support plugins,
> > > but don't support linker scripts this way?
> > 
> > Oops, gold does not support scripts, outputted from plugins :(
> > "error: SECTIONS seen after other input files; try -T/--script"
> > 
> > Probably, we should update default linker scripts in binutils?
> > But without latest ld/gold all binaries compiled without -flto and with
> > offloading will contain intermediate bytecode...
> 
> Actually, this issue is not due to outputting a script from a plugin,
> gold just does not support partial linker scripts:
> https://sourceware.org/bugzilla/show_bug.cgi?id=17451
> 
> So it seems that discarding .gnu.offload_lto_* sections (like it is done for
> .gnu.lto_*) in the default ld and gold scripts is the right way?

I must say I'm not very much familiar with the linker plugin API, but it
surprises me that discarding sections is not something it allows.
Anyway, can you do the partial linker script for the bfd linker (is there
a way to determine from the linker plugin API if it is gold or bfd ld?), and
for gold for the time being perhaps strip the sections in lto-wrapper? and
feed the ET_REL objects with the sections stripped back to the linker
through the plugin API?

	Jakub

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

[Cary Coutant]

The linker already has a --strip-lto-sections option, and it's on by
default. I'll approve a patch that modifies gold to recognize
.gnu.offload_lto.* sections as part of --strip-lto-sections.

Really, though, you should be setting the SHF_EXCLUDE bit on these
sections. Do that and no special-casing will be necessary.

Generating a linker script on the fly to discard these sections is, to
me, rather hacky. There are better ways to do it.

-cary


On Thu, Oct 9, 2014 at 11:53 PM, Jakub Jelinek <jakub@redhat.com> wrote:
> On Fri, Oct 10, 2014 at 12:07:03AM +0400, Ilya Verbin wrote:
>> On 09 Oct 16:07, Ilya Verbin wrote:
>> > > > +      /* By default linker does not discard .gnu.offload_lto_* sections.  */
>> > > > +      const char *linker_script = make_temp_file ("_linker_script.x");
>> > > > +      FILE *stream = fopen (linker_script, "w");
>> > > > +      if (!stream)
>> > > > +       fatal_error ("fopen %s: %m", linker_script);
>> > > > +      fprintf (stream, "SECTIONS { /DISCARD/ : { *("
>> > > > +                      OFFLOAD_SECTION_NAME_PREFIX "*) } }\n");
>> > > > +      fclose (stream);
>> > > > +      printf ("%s\n", linker_script);
>> > > > +
>> > > > +      goto finish;
>> > > > +    }
>> > >
>> > > Does this work with gold?  Are there any other linkers that support plugins,
>> > > but don't support linker scripts this way?
>> >
>> > Oops, gold does not support scripts, outputted from plugins :(
>> > "error: SECTIONS seen after other input files; try -T/--script"
>> >
>> > Probably, we should update default linker scripts in binutils?
>> > But without latest ld/gold all binaries compiled without -flto and with
>> > offloading will contain intermediate bytecode...
>>
>> Actually, this issue is not due to outputting a script from a plugin,
>> gold just does not support partial linker scripts:
>> https://sourceware.org/bugzilla/show_bug.cgi?id=17451
>>
>> So it seems that discarding .gnu.offload_lto_* sections (like it is done for
>> .gnu.lto_*) in the default ld and gold scripts is the right way?
>
> I must say I'm not very much familiar with the linker plugin API, but it
> surprises me that discarding sections is not something it allows.
> Anyway, can you do the partial linker script for the bfd linker (is there
> a way to determine from the linker plugin API if it is gold or bfd ld?), and
> for gold for the time being perhaps strip the sections in lto-wrapper? and
> feed the ET_REL objects with the sections stripped back to the linker
> through the plugin API?
>
>         Jakub

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

[Jakub Jelinek]

On Fri, Oct 10, 2014 at 09:51:02AM -0700, Cary Coutant wrote:
> The linker already has a --strip-lto-sections option, and it's on by
> default. I'll approve a patch that modifies gold to recognize
> .gnu.offload_lto.* sections as part of --strip-lto-sections.
> 
> Really, though, you should be setting the SHF_EXCLUDE bit on these
> sections. Do that and no special-casing will be necessary.

For that I guess
lhd_begin_section
would need to replace:
  section = get_section (name, SECTION_DEBUG, NULL);
with:
  section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL);
either just for the .gnu.offload_lto prefixed section, or all.
The question is what will old assemblers and/or linkers do with that, and
if there are any that support linker plugins, but not SHF_EXCLUDE.

	Jakub

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

[Cary Coutant]

> The question is what will old assemblers and/or linkers do with that, and
> if there are any that support linker plugins, but not SHF_EXCLUDE.

If it helps answer that question, SHF_EXCLUDE support has been in gold
for 6 years, and in gas for 4.

-cary

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

[Ilya Verbin]

On 10 Oct 19:00, Jakub Jelinek wrote:
> On Fri, Oct 10, 2014 at 09:51:02AM -0700, Cary Coutant wrote:
> > The linker already has a --strip-lto-sections option, and it's on by
> > default. I'll approve a patch that modifies gold to recognize
> > .gnu.offload_lto.* sections as part of --strip-lto-sections.
> > 
> > Really, though, you should be setting the SHF_EXCLUDE bit on these
> > sections. Do that and no special-casing will be necessary.
> 
> For that I guess
> lhd_begin_section
> would need to replace:
>   section = get_section (name, SECTION_DEBUG, NULL);
> with:
>   section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL);
> either just for the .gnu.offload_lto prefixed section, or all.
> The question is what will old assemblers and/or linkers do with that, and
> if there are any that support linker plugins, but not SHF_EXCLUDE.

I've tried to set SECTION_EXCLUDE bit with as+ld version 2.20.51 and got a lot
of warnings like:

/tmp/ccg7P7iS.s:2: Warning: entity size for SHF_MERGE not specified
/tmp/ccg7P7iS.s:2: Warning: group name for SHF_GROUP not specified
as: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166'
/usr/bin/ld: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166'

I think, it can be placed under such ifdef:

#if defined (HAVE_SECTION_EXCLUDE) && HAVE_SECTION_EXCLUDE == 1
  section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL);
#else
  section = get_section (name, SECTION_DEBUG, NULL);
#endif

Currently there is HAVE_GAS_SECTION_EXCLUDE implemented in gcc/configure.ac, and
HAVE_SECTION_EXCLUDE can use it + check a version of the linker.


As for old assemblers and/or linkers, which doesn't support SHF_EXCLUDE:

On 10 Oct 08:53, Jakub Jelinek wrote:
> Anyway, can you do the partial linker script for the bfd linker (is there
> a way to determine from the linker plugin API if it is gold or bfd ld?), and
> for gold for the time being perhaps strip the sections in lto-wrapper? and
> feed the ET_REL objects with the sections stripped back to the linker
> through the plugin API?

Yes, it's possible to determine the versions of gold and ld.bfd from the plugin
(plugin-api.h:ld_plugin_tag).  But the problem is not only with plugins.
If linker plugin in disabled, collect2 runs lto-wrapper directly to compile
offload_lto sections and/or to perform usual LTO:
https://gcc.gnu.org/wiki/Offloading#Compilation_without_-flto_and_without_linker_plugin
And if -flto is absent, but offload_lto sections were emitted, the final binary
will contain them.  Perhaps, collect2 could execute and parse `ld --version` to
determine whether it supports SHF_EXCLUDE.

Then, we need somehow pass this information from plugin or collect2 to
lto-wrapper.
Maybe using an option like --strip-offload-sections, but currently lto-wrapper
doesn't have its own options.  And finally, under this option lto-wrapper will
execute `objcopy --remove-section=<offload_lto>` and return stripped objects
back to the linker or to collect2.

So, should I implement this stuff, or maybe just leave offload_lto sections in
the case of old binutils? :)

Thanks,
  -- Ilya

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

[Jakub Jelinek]

On Tue, Oct 14, 2014 at 02:42:47AM +0400, Ilya Verbin wrote:
> > For that I guess
> > lhd_begin_section
> > would need to replace:
> >   section = get_section (name, SECTION_DEBUG, NULL);
> > with:
> >   section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL);
> > either just for the .gnu.offload_lto prefixed section, or all.
> > The question is what will old assemblers and/or linkers do with that, and
> > if there are any that support linker plugins, but not SHF_EXCLUDE.
> 
> I've tried to set SECTION_EXCLUDE bit with as+ld version 2.20.51 and got a lot
> of warnings like:
> 
> /tmp/ccg7P7iS.s:2: Warning: entity size for SHF_MERGE not specified
> /tmp/ccg7P7iS.s:2: Warning: group name for SHF_GROUP not specified
> as: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166'
> /usr/bin/ld: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166'
> 
> I think, it can be placed under such ifdef:
> 
> #if defined (HAVE_SECTION_EXCLUDE) && HAVE_SECTION_EXCLUDE == 1
>   section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL);
> #else
>   section = get_section (name, SECTION_DEBUG, NULL);
> #endif
> 
> Currently there is HAVE_GAS_SECTION_EXCLUDE implemented in gcc/configure.ac, and
> HAVE_SECTION_EXCLUDE can use it + check a version of the linker.

My preference would be to add the | SECTION_EXCLUDE unconditionally, and
instead guard the
  if (flags & SECTION_EXCLUDE)
    *f++ = 'e';
in varasm.c (default_elf_asm_named_section).  The only other user of
SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change
ok with you?

If you have new gas and old linker, I'd expect it would just ignore
SHF_EXCLUDE.

	Jakub

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

[Ilya Verbin]

On 14 Oct 11:40, Jakub Jelinek wrote:
> My preference would be to add the | SECTION_EXCLUDE unconditionally, and
> instead guard the
>   if (flags & SECTION_EXCLUDE)
>     *f++ = 'e';
> in varasm.c (default_elf_asm_named_section).  The only other user of
> SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change
> ok with you?

Done.  But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac:

gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],,
  [.section foo1,"e"
  .byte 0,0,0,0])

does not work properly.  Maybe it works on "cygwin* | pe | mingw32* | interix*"
targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude
sections) it successfully assembles conftest.s into conftest.o (with warnings)
and HAVE_GAS_SECTION_EXCLUDE becomes defined.

$ cat conftest.s
.section foo1,"e"
.byte 0,0,0,0

$ as conftest.s -o conftest.o
conftest.s: Assembler messages:
conftest.s:1: Warning: entity size for SHF_MERGE not specified
conftest.s:1: Warning: group name for SHF_GROUP not specified
as: a.out: warning: sh_link not set for section `foo1'

$ readelf -t conftest.o
[4] foo1
PROGBITS         0000000000000000  0000000000000040  0
0000000000000004 0000000000000000  0                 1
[fffffffffffffdef]: WRITE, ALLOC, EXEC, STRINGS, INFO LINK, LINK ORDER, OS NONCONF, TLS, ORDERED, EXCLUDE, OS (000000000ff00000), PROC (0000000030000000), UNKNOWN (ffffffff000ff808)

And therefore instead of assembling "e" flag, I had to check the version of gas.
If it is >= 2.21.51 then HAVE_GAS_SECTION_EXCLUDE becomes defined.

> If you have new gas and old linker, I'd expect it would just ignore
> SHF_EXCLUDE.

Yes, gas 2.21.51 + ld 2.20.51 build sections with SHF_EXCLUDE bit without
warnings, ld just ignores it.

I've updated the patch.  It also fixes the issue with paths to accel specs.

On 09 Oct 14:08, Bernd Schmidt wrote:
> >>Why do you want to disable specs reading for the accel compiler?
> >>Then users won't have the possibility to override defaults etc. easily...
> I suspect the paths aren't right. If that can be fixed it should be
> fine to remove the ifdef.

Bootstrapped and regtested on i686-linux and x86_64-linux.  Is it OK now?

Thanks,
  -- Ilya


gcc/
	* configure: Regenerate.
	* configure.ac (HAVE_GAS_SECTION_EXCLUDE): Define if GNU assembler
	version >= 2.21.51 is used.
	* gcc.c (spec_host_machine, accel_dir_suffix): New variables.
	(process_command): Tweak path construction for the possibility
	of being configured as an offload compiler.
	(driver::maybe_putenv_OFFLOAD_TARGETS): New function.
	(driver::main): Call maybe_putenv_OFFLOAD_TARGETS.
	(driver::set_up_specs): Tweak path construction for the possibility of
	being configured as an offload compiler.
	* langhooks.c (lhd_begin_section): Set SECTION_EXCLUDE flag.
	* lto-wrapper.c (OFFLOAD_TARGET_NAMES_ENV): Define.
	(offload_names, offloadbegin, offloadend): New static variables.
	(free_array_of_ptrs, parse_env_var, access_check, compile_offload_image)
	(compile_images_for_offload_targets, copy_file, find_offloadbeginend):
	New static functions.
	(run_gcc): Determine whether offload sections are present.  If so, run
	compile_images_for_offload_targets and return the names of new generated
	objects to linker.  If there are offload sections, but no LTO sections,
	then return the copies of input objects without link-time recompilation.
	* varasm.c (default_elf_asm_named_section): Guard SECTION_EXCLUDE with
	ifdef HAVE_GAS_SECTION_EXCLUDE.
lto-plugin/
	* lto-plugin.c (OFFLOAD_SECTION, OFFLOAD_SECTION_LEN): Define.
	(struct plugin_objfile): Add new field "offload".
	(process_offload_section): New static function.
	(claim_file_handler): Claim file if it contains offload sections.

---

diff --git a/gcc/configure b/gcc/configure
index ff1e398..5e61ac8 100755
--- a/gcc/configure
+++ b/gcc/configure
@@ -24755,13 +24755,37 @@ fi
 { $as_echo "$as_me:${as_lineno-$LINENO}: result: $gcc_cv_as_section_has_e" >&5
 $as_echo "$gcc_cv_as_section_has_e" >&6; }
 
+	;;
+    esac
+
+    # GAS versions before 2.21.51 do not support the section flag 'e' for
+    # specifying an excluded section.
+    if test "$gcc_cv_as_section_has_e" != "yes"; then
+      as_ver=`$gcc_cv_as --version 2>/dev/null | sed 1q`
+      if echo "$as_ver" | grep GNU > /dev/null; then
+	as_vers=`echo $as_ver | sed -n \
+	  -e 's,^.*[	 ]\([0-9][0-9]*\.[0-9][0-9]*.*\)$,\1,p'`
+	as_major=`expr "$as_vers" : '\([0-9]*\)'`
+        as_minor=`expr "$as_vers" : '[0-9]*\.\([0-9]*\)'`
+	as_patch=`expr "$as_vers" : '[0-9]*\.[0-9]*\.\([0-9]*\)'`
+        if test $as_major -ge 3; then
+	  gcc_cv_as_section_has_e=yes
+	elif test $as_major -eq 2; then
+	  if test $as_minor -ge 22; then
+	    gcc_cv_as_section_has_e=yes
+	  elif test $as_minor -eq 21; then
+	    if test $as_patch -ge 51; then
+	      gcc_cv_as_section_has_e=yes
+	    fi
+	  fi
+        fi
+      fi
+    fi
 
 cat >>confdefs.h <<_ACEOF
 #define HAVE_GAS_SECTION_EXCLUDE `if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`
 _ACEOF
 
-	;;
-    esac
 
     { $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for filds and fists mnemonics" >&5
 $as_echo_n "checking assembler for filds and fists mnemonics... " >&6; }
diff --git a/gcc/configure.ac b/gcc/configure.ac
index 05a55f4..5edbebe 100644
--- a/gcc/configure.ac
+++ b/gcc/configure.ac
@@ -3843,12 +3843,36 @@ foo:	nop
 	  [2,22,51],,
 [.section foo1,"e"
 .byte 0,0,0,0])
-	AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE,
-	  [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`],
-          [Define if your assembler supports specifying the section flag e.])
 	;;
     esac
 
+    # GAS versions before 2.21.51 do not support the section flag 'e' for
+    # specifying an excluded section.
+    [if test "$gcc_cv_as_section_has_e" != "yes"; then
+      as_ver=`$gcc_cv_as --version 2>/dev/null | sed 1q`
+      if echo "$as_ver" | grep GNU > /dev/null; then
+	as_vers=`echo $as_ver | sed -n \
+	  -e 's,^.*[	 ]\([0-9][0-9]*\.[0-9][0-9]*.*\)$,\1,p'`
+	as_major=`expr "$as_vers" : '\([0-9]*\)'`
+        as_minor=`expr "$as_vers" : '[0-9]*\.\([0-9]*\)'`
+	as_patch=`expr "$as_vers" : '[0-9]*\.[0-9]*\.\([0-9]*\)'`
+        if test $as_major -ge 3; then
+	  gcc_cv_as_section_has_e=yes
+	elif test $as_major -eq 2; then
+	  if test $as_minor -ge 22; then
+	    gcc_cv_as_section_has_e=yes
+	  elif test $as_minor -eq 21; then
+	    if test $as_patch -ge 51; then
+	      gcc_cv_as_section_has_e=yes
+	    fi
+	  fi
+        fi
+      fi
+    fi]
+    AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE,
+      [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`],
+      [Define if your assembler supports specifying the section flag e.])
+
     gcc_GAS_CHECK_FEATURE([filds and fists mnemonics],
        gcc_cv_as_ix86_filds,,,
        [filds mem; fists mem],,
diff --git a/gcc/gcc.c b/gcc/gcc.c
index 71c76f8..4619fe7 100644
--- a/gcc/gcc.c
+++ b/gcc/gcc.c
@@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION;
 /* The target machine.  */
 
 static const char *spec_machine = DEFAULT_TARGET_MACHINE;
+static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE;
 
 /* Nonzero if cross-compiling.
    When -b is used, the value comes from the `specs' file.  */
@@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2
    relative to the driver.  */
 static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX;
 
+/* A prefix to be used when this is an accelerator compiler.  */
+static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX;
+
 /* Subdirectory to use for locating libraries.  Set by
    set_multilib_dir based on the compilation options.  */
 
@@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count,
     }
 
   gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
-  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
+  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
 			    dir_separator_str, NULL);
 
   /* Look for tools relative to the location from which the driver is
      running, or, if that is not available, the configured prefix.  */
   tooldir_prefix
     = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix,
-	      spec_machine, dir_separator_str,
-	      spec_version, dir_separator_str, tooldir_prefix2, NULL);
+	      spec_host_machine, dir_separator_str, spec_version,
+	      accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL);
   free (tooldir_prefix2);
 
   add_prefix (&exec_prefixes,
@@ -6742,6 +6746,7 @@ class driver
   void set_up_specs () const;
   void putenv_COLLECT_GCC (const char *argv0) const;
   void maybe_putenv_COLLECT_LTO_WRAPPER () const;
+  void maybe_putenv_OFFLOAD_TARGETS () const;
   void handle_unrecognized_options () const;
   int maybe_print_and_exit () const;
   bool prepare_infiles ();
@@ -6784,6 +6789,7 @@ driver::main (int argc, char **argv)
   set_up_specs ();
   putenv_COLLECT_GCC (argv[0]);
   maybe_putenv_COLLECT_LTO_WRAPPER ();
+  maybe_putenv_OFFLOAD_TARGETS ();
   handle_unrecognized_options ();
 
   if (!maybe_print_and_exit ())
@@ -6953,6 +6959,7 @@ driver::build_multilib_strings () const
 void
 driver::set_up_specs () const
 {
+  const char *spec_machine_suffix;
   char *specs_file;
   size_t i;
 
@@ -6976,8 +6983,8 @@ driver::set_up_specs () const
 
   /* Read specs from a file if there is one.  */
 
-  machine_suffix = concat (spec_machine, dir_separator_str,
-			   spec_version, dir_separator_str, NULL);
+  machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version,
+			   accel_dir_suffix, dir_separator_str, NULL);
   just_machine_suffix = concat (spec_machine, dir_separator_str, NULL);
 
   specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true);
@@ -6987,13 +6994,18 @@ driver::set_up_specs () const
   else
     init_spec ();
 
-  /* We need to check standard_exec_prefix/just_machine_suffix/specs
+#ifdef ACCEL_COMPILER
+  spec_machine_suffix = machine_suffix;
+#else
+  spec_machine_suffix = just_machine_suffix;
+#endif
+
+  /* We need to check standard_exec_prefix/spec_machine_suffix/specs
      for any override of as, ld and libraries.  */
   specs_file = (char *) alloca (strlen (standard_exec_prefix)
-		       + strlen (just_machine_suffix) + sizeof ("specs"));
-
+		       + strlen (spec_machine_suffix) + sizeof ("specs"));
   strcpy (specs_file, standard_exec_prefix);
-  strcat (specs_file, just_machine_suffix);
+  strcat (specs_file, spec_machine_suffix);
   strcat (specs_file, "specs");
   if (access (specs_file, R_OK) == 0)
     read_specs (specs_file, true, false);
@@ -7175,8 +7187,9 @@ driver::set_up_specs () const
 
   /* If we have a GCC_EXEC_PREFIX envvar, modify it for cpp's sake.  */
   if (gcc_exec_prefix)
-    gcc_exec_prefix = concat (gcc_exec_prefix, spec_machine, dir_separator_str,
-			      spec_version, dir_separator_str, NULL);
+    gcc_exec_prefix = concat (gcc_exec_prefix, spec_host_machine,
+			      dir_separator_str, spec_version,
+			      accel_dir_suffix, dir_separator_str, NULL);
 
   /* Now we have the specs.
      Set the `valid' bits for switches that match anything in any spec.  */
@@ -7227,6 +7240,21 @@ driver::maybe_putenv_COLLECT_LTO_WRAPPER () const
 
 }
 
+/* Set up to remember the names of offload targets.  */
+
+void
+driver::maybe_putenv_OFFLOAD_TARGETS () const
+{
+  if (strlen (OFFLOAD_TARGETS) > 0)
+    {
+      obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=",
+		    sizeof ("OFFLOAD_TARGET_NAMES=") - 1);
+      obstack_grow (&collect_obstack, OFFLOAD_TARGETS,
+		    strlen (OFFLOAD_TARGETS) + 1);
+      xputenv (XOBFINISH (&collect_obstack, char *));
+    }
+}
+
 /* Reject switches that no pass was interested in.  */
 
 void
diff --git a/gcc/langhooks.c b/gcc/langhooks.c
index 7d4c294..4bdeaa0 100644
--- a/gcc/langhooks.c
+++ b/gcc/langhooks.c
@@ -660,7 +660,7 @@ lhd_begin_section (const char *name)
     saved_section = text_section;
 
   /* Create a new section and switch to it.  */
-  section = get_section (name, SECTION_DEBUG, NULL);
+  section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL);
   switch_to_section (section);
 }
 
diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c
index 8033b15..cbda36b 100644
--- a/gcc/lto-wrapper.c
+++ b/gcc/lto-wrapper.c
@@ -49,6 +49,10 @@ along with GCC; see the file COPYING3.  If not see
 #include "lto-section-names.h"
 #include "collect-utils.h"
 
+/* Environment variable, used for passing the names of offload targets from GCC
+   driver to lto-wrapper.  */
+#define OFFLOAD_TARGET_NAMES_ENV	"OFFLOAD_TARGET_NAMES"
+
 enum lto_mode_d {
   LTO_MODE_NONE,			/* Not doing LTO.  */
   LTO_MODE_LTO,				/* Normal LTO.  */
@@ -63,6 +67,8 @@ static char *flto_out;
 static unsigned int nr;
 static char **input_names;
 static char **output_names;
+static char **offload_names;
+static const char *offloadbegin, *offloadend;
 static char *makefile;
 
 const char tool_name[] = "lto-wrapper";
@@ -364,6 +370,223 @@ merge_and_complain (struct cl_decoded_option **decoded_options,
     }
 }
 
+/* Auxiliary function that frees elements of PTR and PTR itself.
+   N is number of elements to be freed.  If PTR is NULL, nothing is freed.
+   If an element is NULL, subsequent elements are not freed.  */
+
+static void **
+free_array_of_ptrs (void **ptr, unsigned n)
+{
+  if (!ptr)
+    return NULL;
+  for (unsigned i = 0; i < n; i++)
+    {
+      if (!ptr[i])
+	break;
+      free (ptr[i]);
+    }
+  free (ptr);
+  return NULL;
+}
+
+/* Parse STR, saving found tokens into PVALUES and return their number.
+   Tokens are assumed to be delimited by ':'.  If APPEND is non-null,
+   append it to every token we find.  */
+
+static unsigned
+parse_env_var (const char *str, char ***pvalues, const char *append)
+{
+  const char *curval, *nextval;
+  char **values;
+  unsigned num = 1, i;
+
+  curval = strchr (str, ':');
+  while (curval)
+    {
+      num++;
+      curval = strchr (curval + 1, ':');
+    }
+
+  values = (char**) xmalloc (num * sizeof (char*));
+  curval = str;
+  nextval = strchrnul (curval, ':');
+
+  int append_len = append ? strlen (append) : 0;
+  for (i = 0; i < num; i++)
+    {
+      int l = nextval - curval;
+      values[i] = (char*) xmalloc (l + 1 + append_len);
+      memcpy (values[i], curval, l);
+      values[i][l] = 0;
+      if (append)
+	strcat (values[i], append);
+      curval = nextval + 1;
+      nextval = strchrnul (curval, ':');
+    }
+  *pvalues = values;
+  return num;
+}
+
+/* Check whether NAME can be accessed in MODE.  This is like access,
+   except that it never considers directories to be executable.  */
+
+static int
+access_check (const char *name, int mode)
+{
+  if (mode == X_OK)
+    {
+      struct stat st;
+
+      if (stat (name, &st) < 0
+	  || S_ISDIR (st.st_mode))
+	return -1;
+    }
+
+  return access (name, mode);
+}
+
+/* Prepare a target image for offload TARGET, using mkoffload tool from
+   COMPILER_PATH.  Return the name of the resultant object file.  */
+
+static char *
+compile_offload_image (const char *target, const char *compiler_path,
+		       unsigned in_argc, char *in_argv[])
+{
+  char *filename = NULL;
+  char **argv;
+  char *suffix
+    = XALLOCAVEC (char, sizeof ("/accel//mkoffload") + strlen (target));
+  strcpy (suffix, "/accel/");
+  strcat (suffix, target);
+  strcat (suffix, "/mkoffload");
+
+  char **paths = NULL;
+  unsigned n_paths = parse_env_var (compiler_path, &paths, suffix);
+
+  const char *compiler = NULL;
+  for (unsigned i = 0; i < n_paths; i++)
+    if (access_check (paths[i], X_OK) == 0)
+      {
+	compiler = paths[i];
+	break;
+      }
+
+  if (compiler)
+    {
+      /* Generate temporary output file name.  */
+      filename = make_temp_file (".target.o");
+
+      struct obstack argv_obstack;
+      obstack_init (&argv_obstack);
+      obstack_ptr_grow (&argv_obstack, compiler);
+      obstack_ptr_grow (&argv_obstack, "-o");
+      obstack_ptr_grow (&argv_obstack, filename);
+
+      for (unsigned i = 1; i < in_argc; i++)
+	obstack_ptr_grow (&argv_obstack, in_argv[i]);
+      obstack_ptr_grow (&argv_obstack, NULL);
+
+      argv = XOBFINISH (&argv_obstack, char **);
+      fork_execute (argv[0], argv, true);
+      obstack_free (&argv_obstack, NULL);
+    }
+
+  free_array_of_ptrs ((void **) paths, n_paths);
+  return filename;
+}
+
+
+/* The main routine dealing with offloading.
+   The routine builds a target image for each offload target.  IN_ARGC and
+   IN_ARGV specify options and input object files.  As all of them could contain
+   target sections, we pass them all to target compilers.  */
+
+static void
+compile_images_for_offload_targets (unsigned in_argc, char *in_argv[])
+{
+  char **names = NULL;
+  const char *target_names = getenv (OFFLOAD_TARGET_NAMES_ENV);
+  if (!target_names)
+    return;
+  unsigned num_targets = parse_env_var (target_names, &names, NULL);
+
+  const char *compiler_path = getenv ("COMPILER_PATH");
+  if (!compiler_path)
+    goto out;
+
+  /* Prepare an image for each target and save the name of the resultant object
+     file to the OFFLOAD_NAMES array.  It is terminated by a NULL entry.  */
+  offload_names = XCNEWVEC (char *, num_targets + 1);
+  for (unsigned i = 0; i < num_targets; i++)
+    {
+      offload_names[i] = compile_offload_image (names[i], compiler_path,
+						in_argc, in_argv);
+      if (!offload_names[i])
+	fatal_error ("problem with building target image for %s\n", names[i]);
+    }
+
+ out:
+  free_array_of_ptrs ((void **) names, num_targets);
+}
+
+/* Copy a file from SRC to DEST.  */
+
+static void
+copy_file (const char *dest, const char *src)
+{
+  FILE *d = fopen (dest, "wb");
+  FILE *s = fopen (src, "rb");
+  char buffer[512];
+  while (!feof (s))
+    {
+      size_t len = fread (buffer, 1, 512, s);
+      if (ferror (s) != 0)
+	fatal_error ("reading input file");
+      if (len > 0)
+	{
+	  fwrite (buffer, 1, len, d);
+	  if (ferror (d) != 0)
+	    fatal_error ("writing output file");
+	}
+    }
+}
+
+/* Find the crtoffloadbegin.o and crtoffloadend.o files in LIBRARY_PATH, make
+   copies and store the names of the copies in offloadbegin and offloadend.  */
+
+static void
+find_offloadbeginend (void)
+{
+  char **paths = NULL;
+  const char *library_path = getenv ("LIBRARY_PATH");
+  if (!library_path)
+    return;
+  unsigned n_paths = parse_env_var (library_path, &paths, "/crtoffloadbegin.o");
+
+  unsigned i;
+  for (i = 0; i < n_paths; i++)
+    if (access_check (paths[i], R_OK) == 0)
+      {
+	size_t len = strlen (paths[i]);
+	char *tmp = xstrdup (paths[i]);
+	strcpy (paths[i] + len - strlen ("begin.o"), "end.o");
+	if (access_check (paths[i], R_OK) != 0)
+	  fatal_error ("installation error, can't find crtoffloadend.o");
+	/* The linker will delete the filenames we give it, so make
+	   copies.  */
+	offloadbegin = make_temp_file (".o");
+	offloadend = make_temp_file (".o");
+	copy_file (offloadbegin, tmp);
+	copy_file (offloadend, paths[i]);
+	free (tmp);
+	break;
+      }
+  if (i == n_paths)
+    fatal_error ("installation error, can't find crtoffloadbegin.o");
+
+  free_array_of_ptrs ((void **) paths, n_paths);
+}
+
 /* Execute gcc. ARGC is the number of arguments. ARGV contains the arguments. */
 
 static void
@@ -384,6 +607,8 @@ run_gcc (unsigned argc, char *argv[])
   unsigned int decoded_options_count;
   struct obstack argv_obstack;
   int new_head_argc;
+  bool have_lto = false;
+  bool have_offload = false;
 
   /* Get the driver and options.  */
   collect_gcc = getenv ("COLLECT_GCC");
@@ -432,6 +657,9 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
+      if (simple_object_find_section (sobj, OFFLOAD_SECTION_NAME_PREFIX ".opts",
+				      &offset, &length, &errmsg, &err))
+	have_offload = true;
       if (!simple_object_find_section (sobj, LTO_SECTION_NAME_PREFIX "." "opts",
 				       &offset, &length, &errmsg, &err))
 	{
@@ -439,6 +667,7 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
+      have_lto = true;
       lseek (fd, file_offset + offset, SEEK_SET);
       data = (char *)xmalloc (length);
       read (fd, data, length);
@@ -633,6 +862,43 @@ run_gcc (unsigned argc, char *argv[])
   /* Remember at which point we can scrub args to re-use the commons.  */
   new_head_argc = obstack_object_size (&argv_obstack) / sizeof (void *);
 
+  if (have_offload)
+    {
+      compile_images_for_offload_targets (argc, argv);
+      if (offload_names)
+	{
+	  find_offloadbeginend ();
+	  for (i = 0; offload_names[i]; i++)
+	    printf ("%s\n", offload_names[i]);
+	  free_array_of_ptrs ((void **) offload_names, i);
+	}
+    }
+
+  if (offloadbegin)
+    printf ("%s\n", offloadbegin);
+
+  /* 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 (strncmp (argv[i], "-fresolution=", sizeof ("-fresolution=") - 1))
+	  {
+	    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");
@@ -859,6 +1125,10 @@ cont:
       obstack_free (&env_obstack, NULL);
     }
 
+ finish:
+  if (offloadend)
+    printf ("%s\n", offloadend);
+
   obstack_free (&argv_obstack, NULL);
 }
 
diff --git a/gcc/varasm.c b/gcc/varasm.c
index abb743b..4ae9d58 100644
--- a/gcc/varasm.c
+++ b/gcc/varasm.c
@@ -6141,8 +6141,10 @@ default_elf_asm_named_section (const char *name, unsigned int flags,
 
   if (!(flags & SECTION_DEBUG))
     *f++ = 'a';
+#if defined (HAVE_GAS_SECTION_EXCLUDE) && HAVE_GAS_SECTION_EXCLUDE == 1
   if (flags & SECTION_EXCLUDE)
     *f++ = 'e';
+#endif
   if (flags & SECTION_WRITE)
     *f++ = 'w';
   if (flags & SECTION_CODE)
diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c
index 910e23c..fb6555d 100644
--- a/lto-plugin/lto-plugin.c
+++ b/lto-plugin/lto-plugin.c
@@ -86,6 +86,8 @@ along with this program; see the file COPYING3.  If not see
 
 #define LTO_SECTION_PREFIX	".gnu.lto_.symtab"
 #define LTO_SECTION_PREFIX_LEN	(sizeof (LTO_SECTION_PREFIX) - 1)
+#define OFFLOAD_SECTION		".gnu.offload_lto_.opts"
+#define OFFLOAD_SECTION_LEN	(sizeof (OFFLOAD_SECTION) - 1)
 
 /* The part of the symbol table the plugin has to keep track of. Note that we
    must keep SYMS until all_symbols_read is called to give the linker time to
@@ -111,6 +113,7 @@ struct plugin_symtab
 struct plugin_objfile
 {
   int found;
+  int offload;
   simple_object_read *objfile;
   struct plugin_symtab *out;
   const struct ld_plugin_input_file *file;
@@ -862,6 +865,21 @@ err:
   return 0;
 }
 
+/* Find an offload section of an object file.  */
+
+static int
+process_offload_section (void *data, const char *name, off_t offset, off_t len)
+{
+  if (!strncmp (name, OFFLOAD_SECTION, OFFLOAD_SECTION_LEN))
+    {
+      struct plugin_objfile *obj = (struct plugin_objfile *) data;
+      obj->offload = 1;
+      return 0;
+    }
+
+  return 1;
+}
+
 /* Callback used by gold to check if the plugin will claim FILE. Writes
    the result in CLAIMED. */
 
@@ -899,6 +917,7 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed)
   *claimed = 0;
   obj.file = file;
   obj.found = 0;
+  obj.offload = 0;
   obj.out = &lto_file.symtab;
   errmsg = NULL;
   obj.objfile = simple_object_start_read (file->fd, file->offset, LTO_SEGMENT_NAME,
@@ -920,7 +939,11 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed)
       goto err;
     }
 
-  if (obj.found == 0)
+  if (obj.objfile)
+    simple_object_find_sections (obj.objfile, process_offload_section,
+				 &obj, &err);
+
+  if (obj.found == 0 && obj.offload == 0)
     goto err;
 
   if (obj.found > 1)
-- 
1.7.1

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

[Jakub Jelinek]

On Wed, Oct 15, 2014 at 06:18:56PM +0400, Ilya Verbin wrote:
> On 14 Oct 11:40, Jakub Jelinek wrote:
> > My preference would be to add the | SECTION_EXCLUDE unconditionally, and
> > instead guard the
> >   if (flags & SECTION_EXCLUDE)
> >     *f++ = 'e';
> > in varasm.c (default_elf_asm_named_section).  The only other user of
> > SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change
> > ok with you?
> 
> Done.  But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac:
> 
> gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],,
>   [.section foo1,"e"
>   .byte 0,0,0,0])
> 
> does not work properly.  Maybe it works on "cygwin* | pe | mingw32* | interix*"
> targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude
> sections) it successfully assembles conftest.s into conftest.o (with warnings)
> and HAVE_GAS_SECTION_EXCLUDE becomes defined.

IMHO a version check is wrong (except when using in-tree gas).
I'd suggest just to use [--fatal-warnings] as the 4th argument to
gcc_GAS_CHECK_FEATURE feature, after all, that is what e.g.
gcc_cv_as_shf_merge testing already uses.

	Jakub

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

[Cary Coutant]

> My preference would be to add the | SECTION_EXCLUDE unconditionally, and
> instead guard the
>   if (flags & SECTION_EXCLUDE)
>     *f++ = 'e';
> in varasm.c (default_elf_asm_named_section).  The only other user of
> SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change
> ok with you?

Yes, that sounds fine.

> If you have new gas and old linker, I'd expect it would just ignore
> SHF_EXCLUDE.

Agreed.

-cary

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

[Ilya Verbin]

On 15 Oct 16:38, Jakub Jelinek wrote:
> > Done.  But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac:
> > 
> > gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],,
> >   [.section foo1,"e"
> >   .byte 0,0,0,0])
> > 
> > does not work properly.  Maybe it works on "cygwin* | pe | mingw32* | interix*"
> > targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude
> > sections) it successfully assembles conftest.s into conftest.o (with warnings)
> > and HAVE_GAS_SECTION_EXCLUDE becomes defined.
> 
> IMHO a version check is wrong (except when using in-tree gas).
> I'd suggest just to use [--fatal-warnings] as the 4th argument to
> gcc_GAS_CHECK_FEATURE feature, after all, that is what e.g.
> gcc_cv_as_shf_merge testing already uses.

Fixed.  Patch is updated and retested.

Thanks,
  -- Ilya


gcc/
	* configure: Regenerate.
	* configure.ac: Move the test for section attribute specifier "e" in GAS
	out to all i[34567]86-*-* | x86_64-*-* targets and add --fatal-warnings.
	* gcc.c (spec_host_machine, accel_dir_suffix): New variables.
	(process_command): Tweak path construction for the possibility
	of being configured as an offload compiler.
	(driver::maybe_putenv_OFFLOAD_TARGETS): New function.
	(driver::main): Call maybe_putenv_OFFLOAD_TARGETS.
	(driver::set_up_specs): Tweak path construction for the possibility of
	being configured as an offload compiler.
	* langhooks.c (lhd_begin_section): Set SECTION_EXCLUDE flag.
	* lto-wrapper.c (OFFLOAD_TARGET_NAMES_ENV): Define.
	(offload_names, offloadbegin, offloadend): New static variables.
	(free_array_of_ptrs, parse_env_var, access_check, compile_offload_image)
	(compile_images_for_offload_targets, copy_file, find_offloadbeginend):
	New static functions.
	(run_gcc): Determine whether offload sections are present.  If so, run
	compile_images_for_offload_targets and return the names of new generated
	objects to linker.  If there are offload sections, but no LTO sections,
	then return the copies of input objects without link-time recompilation.
	* varasm.c (default_elf_asm_named_section): Guard SECTION_EXCLUDE with
	ifdef HAVE_GAS_SECTION_EXCLUDE.
lto-plugin/
	* lto-plugin.c (OFFLOAD_SECTION, OFFLOAD_SECTION_LEN): Define.
	(struct plugin_objfile): Add new field "offload".
	(process_offload_section): New static function.
	(claim_file_handler): Claim file if it contains offload sections.

---

diff --git a/gcc/configure b/gcc/configure
index ff1e398..4ef208c 100755
--- a/gcc/configure
+++ b/gcc/configure
@@ -24722,9 +24722,12 @@ $as_echo "$as_me: WARNING: LTO for $target requires binutils >= 2.20.1, but vers
 	      ;;
 	  esac
 	fi
-	# Test if the assembler supports the section flag 'e' for specifying
-	# an excluded section.
-	{ $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for .section with e" >&5
+	;;
+    esac
+
+    # Test if the assembler supports the section flag 'e' for specifying
+    # an excluded section.
+    { $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for .section with e" >&5
 $as_echo_n "checking assembler for .section with e... " >&6; }
 if test "${gcc_cv_as_section_has_e+set}" = set; then :
   $as_echo_n "(cached) " >&6
@@ -24737,7 +24740,7 @@ fi
   elif test x$gcc_cv_as != x; then
     $as_echo '.section foo1,"e"
 .byte 0,0,0,0' > conftest.s
-    if { ac_try='$gcc_cv_as $gcc_cv_as_flags  -o conftest.o conftest.s >&5'
+    if { ac_try='$gcc_cv_as $gcc_cv_as_flags --fatal-warnings -o conftest.o conftest.s >&5'
   { { eval echo "\"\$as_me\":${as_lineno-$LINENO}: \"$ac_try\""; } >&5
   (eval $ac_try) 2>&5
   ac_status=$?
@@ -24760,8 +24763,6 @@ cat >>confdefs.h <<_ACEOF
 #define HAVE_GAS_SECTION_EXCLUDE `if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`
 _ACEOF
 
-	;;
-    esac
 
     { $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for filds and fists mnemonics" >&5
 $as_echo_n "checking assembler for filds and fists mnemonics... " >&6; }
diff --git a/gcc/configure.ac b/gcc/configure.ac
index 05a55f4..0f4bfc6 100644
--- a/gcc/configure.ac
+++ b/gcc/configure.ac
@@ -3837,18 +3837,19 @@ foo:	nop
 	      ;;
 	  esac
 	fi
-	# Test if the assembler supports the section flag 'e' for specifying
-	# an excluded section.
-	gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e,
-	  [2,22,51],,
-[.section foo1,"e"
-.byte 0,0,0,0])
-	AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE,
-	  [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`],
-          [Define if your assembler supports specifying the section flag e.])
 	;;
     esac
 
+    # Test if the assembler supports the section flag 'e' for specifying
+    # an excluded section.
+    gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e,
+      [2,22,51], [--fatal-warnings],
+[.section foo1,"e"
+.byte 0,0,0,0])
+    AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE,
+      [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`],
+      [Define if your assembler supports specifying the section flag e.])
+
     gcc_GAS_CHECK_FEATURE([filds and fists mnemonics],
        gcc_cv_as_ix86_filds,,,
        [filds mem; fists mem],,
diff --git a/gcc/gcc.c b/gcc/gcc.c
index 71c76f8..4619fe7 100644
--- a/gcc/gcc.c
+++ b/gcc/gcc.c
@@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION;
 /* The target machine.  */
 
 static const char *spec_machine = DEFAULT_TARGET_MACHINE;
+static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE;
 
 /* Nonzero if cross-compiling.
    When -b is used, the value comes from the `specs' file.  */
@@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2
    relative to the driver.  */
 static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX;
 
+/* A prefix to be used when this is an accelerator compiler.  */
+static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX;
+
 /* Subdirectory to use for locating libraries.  Set by
    set_multilib_dir based on the compilation options.  */
 
@@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count,
     }
 
   gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
-  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
+  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
 			    dir_separator_str, NULL);
 
   /* Look for tools relative to the location from which the driver is
      running, or, if that is not available, the configured prefix.  */
   tooldir_prefix
     = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix,
-	      spec_machine, dir_separator_str,
-	      spec_version, dir_separator_str, tooldir_prefix2, NULL);
+	      spec_host_machine, dir_separator_str, spec_version,
+	      accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL);
   free (tooldir_prefix2);
 
   add_prefix (&exec_prefixes,
@@ -6742,6 +6746,7 @@ class driver
   void set_up_specs () const;
   void putenv_COLLECT_GCC (const char *argv0) const;
   void maybe_putenv_COLLECT_LTO_WRAPPER () const;
+  void maybe_putenv_OFFLOAD_TARGETS () const;
   void handle_unrecognized_options () const;
   int maybe_print_and_exit () const;
   bool prepare_infiles ();
@@ -6784,6 +6789,7 @@ driver::main (int argc, char **argv)
   set_up_specs ();
   putenv_COLLECT_GCC (argv[0]);
   maybe_putenv_COLLECT_LTO_WRAPPER ();
+  maybe_putenv_OFFLOAD_TARGETS ();
   handle_unrecognized_options ();
 
   if (!maybe_print_and_exit ())
@@ -6953,6 +6959,7 @@ driver::build_multilib_strings () const
 void
 driver::set_up_specs () const
 {
+  const char *spec_machine_suffix;
   char *specs_file;
   size_t i;
 
@@ -6976,8 +6983,8 @@ driver::set_up_specs () const
 
   /* Read specs from a file if there is one.  */
 
-  machine_suffix = concat (spec_machine, dir_separator_str,
-			   spec_version, dir_separator_str, NULL);
+  machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version,
+			   accel_dir_suffix, dir_separator_str, NULL);
   just_machine_suffix = concat (spec_machine, dir_separator_str, NULL);
 
   specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true);
@@ -6987,13 +6994,18 @@ driver::set_up_specs () const
   else
     init_spec ();
 
-  /* We need to check standard_exec_prefix/just_machine_suffix/specs
+#ifdef ACCEL_COMPILER
+  spec_machine_suffix = machine_suffix;
+#else
+  spec_machine_suffix = just_machine_suffix;
+#endif
+
+  /* We need to check standard_exec_prefix/spec_machine_suffix/specs
      for any override of as, ld and libraries.  */
   specs_file = (char *) alloca (strlen (standard_exec_prefix)
-		       + strlen (just_machine_suffix) + sizeof ("specs"));
-
+		       + strlen (spec_machine_suffix) + sizeof ("specs"));
   strcpy (specs_file, standard_exec_prefix);
-  strcat (specs_file, just_machine_suffix);
+  strcat (specs_file, spec_machine_suffix);
   strcat (specs_file, "specs");
   if (access (specs_file, R_OK) == 0)
     read_specs (specs_file, true, false);
@@ -7175,8 +7187,9 @@ driver::set_up_specs () const
 
   /* If we have a GCC_EXEC_PREFIX envvar, modify it for cpp's sake.  */
   if (gcc_exec_prefix)
-    gcc_exec_prefix = concat (gcc_exec_prefix, spec_machine, dir_separator_str,
-			      spec_version, dir_separator_str, NULL);
+    gcc_exec_prefix = concat (gcc_exec_prefix, spec_host_machine,
+			      dir_separator_str, spec_version,
+			      accel_dir_suffix, dir_separator_str, NULL);
 
   /* Now we have the specs.
      Set the `valid' bits for switches that match anything in any spec.  */
@@ -7227,6 +7240,21 @@ driver::maybe_putenv_COLLECT_LTO_WRAPPER () const
 
 }
 
+/* Set up to remember the names of offload targets.  */
+
+void
+driver::maybe_putenv_OFFLOAD_TARGETS () const
+{
+  if (strlen (OFFLOAD_TARGETS) > 0)
+    {
+      obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=",
+		    sizeof ("OFFLOAD_TARGET_NAMES=") - 1);
+      obstack_grow (&collect_obstack, OFFLOAD_TARGETS,
+		    strlen (OFFLOAD_TARGETS) + 1);
+      xputenv (XOBFINISH (&collect_obstack, char *));
+    }
+}
+
 /* Reject switches that no pass was interested in.  */
 
 void
diff --git a/gcc/langhooks.c b/gcc/langhooks.c
index 7d4c294..4bdeaa0 100644
--- a/gcc/langhooks.c
+++ b/gcc/langhooks.c
@@ -660,7 +660,7 @@ lhd_begin_section (const char *name)
     saved_section = text_section;
 
   /* Create a new section and switch to it.  */
-  section = get_section (name, SECTION_DEBUG, NULL);
+  section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL);
   switch_to_section (section);
 }
 
diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c
index 8033b15..cbda36b 100644
--- a/gcc/lto-wrapper.c
+++ b/gcc/lto-wrapper.c
@@ -49,6 +49,10 @@ along with GCC; see the file COPYING3.  If not see
 #include "lto-section-names.h"
 #include "collect-utils.h"
 
+/* Environment variable, used for passing the names of offload targets from GCC
+   driver to lto-wrapper.  */
+#define OFFLOAD_TARGET_NAMES_ENV	"OFFLOAD_TARGET_NAMES"
+
 enum lto_mode_d {
   LTO_MODE_NONE,			/* Not doing LTO.  */
   LTO_MODE_LTO,				/* Normal LTO.  */
@@ -63,6 +67,8 @@ static char *flto_out;
 static unsigned int nr;
 static char **input_names;
 static char **output_names;
+static char **offload_names;
+static const char *offloadbegin, *offloadend;
 static char *makefile;
 
 const char tool_name[] = "lto-wrapper";
@@ -364,6 +370,223 @@ merge_and_complain (struct cl_decoded_option **decoded_options,
     }
 }
 
+/* Auxiliary function that frees elements of PTR and PTR itself.
+   N is number of elements to be freed.  If PTR is NULL, nothing is freed.
+   If an element is NULL, subsequent elements are not freed.  */
+
+static void **
+free_array_of_ptrs (void **ptr, unsigned n)
+{
+  if (!ptr)
+    return NULL;
+  for (unsigned i = 0; i < n; i++)
+    {
+      if (!ptr[i])
+	break;
+      free (ptr[i]);
+    }
+  free (ptr);
+  return NULL;
+}
+
+/* Parse STR, saving found tokens into PVALUES and return their number.
+   Tokens are assumed to be delimited by ':'.  If APPEND is non-null,
+   append it to every token we find.  */
+
+static unsigned
+parse_env_var (const char *str, char ***pvalues, const char *append)
+{
+  const char *curval, *nextval;
+  char **values;
+  unsigned num = 1, i;
+
+  curval = strchr (str, ':');
+  while (curval)
+    {
+      num++;
+      curval = strchr (curval + 1, ':');
+    }
+
+  values = (char**) xmalloc (num * sizeof (char*));
+  curval = str;
+  nextval = strchrnul (curval, ':');
+
+  int append_len = append ? strlen (append) : 0;
+  for (i = 0; i < num; i++)
+    {
+      int l = nextval - curval;
+      values[i] = (char*) xmalloc (l + 1 + append_len);
+      memcpy (values[i], curval, l);
+      values[i][l] = 0;
+      if (append)
+	strcat (values[i], append);
+      curval = nextval + 1;
+      nextval = strchrnul (curval, ':');
+    }
+  *pvalues = values;
+  return num;
+}
+
+/* Check whether NAME can be accessed in MODE.  This is like access,
+   except that it never considers directories to be executable.  */
+
+static int
+access_check (const char *name, int mode)
+{
+  if (mode == X_OK)
+    {
+      struct stat st;
+
+      if (stat (name, &st) < 0
+	  || S_ISDIR (st.st_mode))
+	return -1;
+    }
+
+  return access (name, mode);
+}
+
+/* Prepare a target image for offload TARGET, using mkoffload tool from
+   COMPILER_PATH.  Return the name of the resultant object file.  */
+
+static char *
+compile_offload_image (const char *target, const char *compiler_path,
+		       unsigned in_argc, char *in_argv[])
+{
+  char *filename = NULL;
+  char **argv;
+  char *suffix
+    = XALLOCAVEC (char, sizeof ("/accel//mkoffload") + strlen (target));
+  strcpy (suffix, "/accel/");
+  strcat (suffix, target);
+  strcat (suffix, "/mkoffload");
+
+  char **paths = NULL;
+  unsigned n_paths = parse_env_var (compiler_path, &paths, suffix);
+
+  const char *compiler = NULL;
+  for (unsigned i = 0; i < n_paths; i++)
+    if (access_check (paths[i], X_OK) == 0)
+      {
+	compiler = paths[i];
+	break;
+      }
+
+  if (compiler)
+    {
+      /* Generate temporary output file name.  */
+      filename = make_temp_file (".target.o");
+
+      struct obstack argv_obstack;
+      obstack_init (&argv_obstack);
+      obstack_ptr_grow (&argv_obstack, compiler);
+      obstack_ptr_grow (&argv_obstack, "-o");
+      obstack_ptr_grow (&argv_obstack, filename);
+
+      for (unsigned i = 1; i < in_argc; i++)
+	obstack_ptr_grow (&argv_obstack, in_argv[i]);
+      obstack_ptr_grow (&argv_obstack, NULL);
+
+      argv = XOBFINISH (&argv_obstack, char **);
+      fork_execute (argv[0], argv, true);
+      obstack_free (&argv_obstack, NULL);
+    }
+
+  free_array_of_ptrs ((void **) paths, n_paths);
+  return filename;
+}
+
+
+/* The main routine dealing with offloading.
+   The routine builds a target image for each offload target.  IN_ARGC and
+   IN_ARGV specify options and input object files.  As all of them could contain
+   target sections, we pass them all to target compilers.  */
+
+static void
+compile_images_for_offload_targets (unsigned in_argc, char *in_argv[])
+{
+  char **names = NULL;
+  const char *target_names = getenv (OFFLOAD_TARGET_NAMES_ENV);
+  if (!target_names)
+    return;
+  unsigned num_targets = parse_env_var (target_names, &names, NULL);
+
+  const char *compiler_path = getenv ("COMPILER_PATH");
+  if (!compiler_path)
+    goto out;
+
+  /* Prepare an image for each target and save the name of the resultant object
+     file to the OFFLOAD_NAMES array.  It is terminated by a NULL entry.  */
+  offload_names = XCNEWVEC (char *, num_targets + 1);
+  for (unsigned i = 0; i < num_targets; i++)
+    {
+      offload_names[i] = compile_offload_image (names[i], compiler_path,
+						in_argc, in_argv);
+      if (!offload_names[i])
+	fatal_error ("problem with building target image for %s\n", names[i]);
+    }
+
+ out:
+  free_array_of_ptrs ((void **) names, num_targets);
+}
+
+/* Copy a file from SRC to DEST.  */
+
+static void
+copy_file (const char *dest, const char *src)
+{
+  FILE *d = fopen (dest, "wb");
+  FILE *s = fopen (src, "rb");
+  char buffer[512];
+  while (!feof (s))
+    {
+      size_t len = fread (buffer, 1, 512, s);
+      if (ferror (s) != 0)
+	fatal_error ("reading input file");
+      if (len > 0)
+	{
+	  fwrite (buffer, 1, len, d);
+	  if (ferror (d) != 0)
+	    fatal_error ("writing output file");
+	}
+    }
+}
+
+/* Find the crtoffloadbegin.o and crtoffloadend.o files in LIBRARY_PATH, make
+   copies and store the names of the copies in offloadbegin and offloadend.  */
+
+static void
+find_offloadbeginend (void)
+{
+  char **paths = NULL;
+  const char *library_path = getenv ("LIBRARY_PATH");
+  if (!library_path)
+    return;
+  unsigned n_paths = parse_env_var (library_path, &paths, "/crtoffloadbegin.o");
+
+  unsigned i;
+  for (i = 0; i < n_paths; i++)
+    if (access_check (paths[i], R_OK) == 0)
+      {
+	size_t len = strlen (paths[i]);
+	char *tmp = xstrdup (paths[i]);
+	strcpy (paths[i] + len - strlen ("begin.o"), "end.o");
+	if (access_check (paths[i], R_OK) != 0)
+	  fatal_error ("installation error, can't find crtoffloadend.o");
+	/* The linker will delete the filenames we give it, so make
+	   copies.  */
+	offloadbegin = make_temp_file (".o");
+	offloadend = make_temp_file (".o");
+	copy_file (offloadbegin, tmp);
+	copy_file (offloadend, paths[i]);
+	free (tmp);
+	break;
+      }
+  if (i == n_paths)
+    fatal_error ("installation error, can't find crtoffloadbegin.o");
+
+  free_array_of_ptrs ((void **) paths, n_paths);
+}
+
 /* Execute gcc. ARGC is the number of arguments. ARGV contains the arguments. */
 
 static void
@@ -384,6 +607,8 @@ run_gcc (unsigned argc, char *argv[])
   unsigned int decoded_options_count;
   struct obstack argv_obstack;
   int new_head_argc;
+  bool have_lto = false;
+  bool have_offload = false;
 
   /* Get the driver and options.  */
   collect_gcc = getenv ("COLLECT_GCC");
@@ -432,6 +657,9 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
+      if (simple_object_find_section (sobj, OFFLOAD_SECTION_NAME_PREFIX ".opts",
+				      &offset, &length, &errmsg, &err))
+	have_offload = true;
       if (!simple_object_find_section (sobj, LTO_SECTION_NAME_PREFIX "." "opts",
 				       &offset, &length, &errmsg, &err))
 	{
@@ -439,6 +667,7 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
+      have_lto = true;
       lseek (fd, file_offset + offset, SEEK_SET);
       data = (char *)xmalloc (length);
       read (fd, data, length);
@@ -633,6 +862,43 @@ run_gcc (unsigned argc, char *argv[])
   /* Remember at which point we can scrub args to re-use the commons.  */
   new_head_argc = obstack_object_size (&argv_obstack) / sizeof (void *);
 
+  if (have_offload)
+    {
+      compile_images_for_offload_targets (argc, argv);
+      if (offload_names)
+	{
+	  find_offloadbeginend ();
+	  for (i = 0; offload_names[i]; i++)
+	    printf ("%s\n", offload_names[i]);
+	  free_array_of_ptrs ((void **) offload_names, i);
+	}
+    }
+
+  if (offloadbegin)
+    printf ("%s\n", offloadbegin);
+
+  /* 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 (strncmp (argv[i], "-fresolution=", sizeof ("-fresolution=") - 1))
+	  {
+	    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");
@@ -859,6 +1125,10 @@ cont:
       obstack_free (&env_obstack, NULL);
     }
 
+ finish:
+  if (offloadend)
+    printf ("%s\n", offloadend);
+
   obstack_free (&argv_obstack, NULL);
 }
 
diff --git a/gcc/varasm.c b/gcc/varasm.c
index abb743b..4ae9d58 100644
--- a/gcc/varasm.c
+++ b/gcc/varasm.c
@@ -6141,8 +6141,10 @@ default_elf_asm_named_section (const char *name, unsigned int flags,
 
   if (!(flags & SECTION_DEBUG))
     *f++ = 'a';
+#if defined (HAVE_GAS_SECTION_EXCLUDE) && HAVE_GAS_SECTION_EXCLUDE == 1
   if (flags & SECTION_EXCLUDE)
     *f++ = 'e';
+#endif
   if (flags & SECTION_WRITE)
     *f++ = 'w';
   if (flags & SECTION_CODE)
diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c
index 910e23c..fb6555d 100644
--- a/lto-plugin/lto-plugin.c
+++ b/lto-plugin/lto-plugin.c
@@ -86,6 +86,8 @@ along with this program; see the file COPYING3.  If not see
 
 #define LTO_SECTION_PREFIX	".gnu.lto_.symtab"
 #define LTO_SECTION_PREFIX_LEN	(sizeof (LTO_SECTION_PREFIX) - 1)
+#define OFFLOAD_SECTION		".gnu.offload_lto_.opts"
+#define OFFLOAD_SECTION_LEN	(sizeof (OFFLOAD_SECTION) - 1)
 
 /* The part of the symbol table the plugin has to keep track of. Note that we
    must keep SYMS until all_symbols_read is called to give the linker time to
@@ -111,6 +113,7 @@ struct plugin_symtab
 struct plugin_objfile
 {
   int found;
+  int offload;
   simple_object_read *objfile;
   struct plugin_symtab *out;
   const struct ld_plugin_input_file *file;
@@ -862,6 +865,21 @@ err:
   return 0;
 }
 
+/* Find an offload section of an object file.  */
+
+static int
+process_offload_section (void *data, const char *name, off_t offset, off_t len)
+{
+  if (!strncmp (name, OFFLOAD_SECTION, OFFLOAD_SECTION_LEN))
+    {
+      struct plugin_objfile *obj = (struct plugin_objfile *) data;
+      obj->offload = 1;
+      return 0;
+    }
+
+  return 1;
+}
+
 /* Callback used by gold to check if the plugin will claim FILE. Writes
    the result in CLAIMED. */
 
@@ -899,6 +917,7 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed)
   *claimed = 0;
   obj.file = file;
   obj.found = 0;
+  obj.offload = 0;
   obj.out = &lto_file.symtab;
   errmsg = NULL;
   obj.objfile = simple_object_start_read (file->fd, file->offset, LTO_SEGMENT_NAME,
@@ -920,7 +939,11 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed)
       goto err;
     }
 
-  if (obj.found == 0)
+  if (obj.objfile)
+    simple_object_find_sections (obj.objfile, process_offload_section,
+				 &obj, &err);
+
+  if (obj.found == 0 && obj.offload == 0)
     goto err;
 
   if (obj.found > 1)
-- 
1.7.1

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

[Jakub Jelinek]

On Thu, Oct 16, 2014 at 03:17:36PM +0400, Ilya Verbin wrote:
> On 15 Oct 16:38, Jakub Jelinek wrote:
> > > Done.  But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac:
> > > 
> > > gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],,
> > >   [.section foo1,"e"
> > >   .byte 0,0,0,0])
> > > 
> > > does not work properly.  Maybe it works on "cygwin* | pe | mingw32* | interix*"
> > > targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude
> > > sections) it successfully assembles conftest.s into conftest.o (with warnings)
> > > and HAVE_GAS_SECTION_EXCLUDE becomes defined.
> > 
> > IMHO a version check is wrong (except when using in-tree gas).
> > I'd suggest just to use [--fatal-warnings] as the 4th argument to
> > gcc_GAS_CHECK_FEATURE feature, after all, that is what e.g.
> > gcc_cv_as_shf_merge testing already uses.
> 
> Fixed.  Patch is updated and retested.

Can you please extract the configure{,.ac}, langhooks.c and varasm.c
bits into a separate patch?  That is preapproved for trunk right now, that
isn't dependent on anything else.

The rest LGTM, but please run it through LTO review (Richard/Honza) too.

	Jakub

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

[Kirill Yukhin]

Hello Richard, Jan,
On 16 Oct 13:22, Jakub Jelinek wrote:
> On Thu, Oct 16, 2014 at 03:17:36PM +0400, Ilya Verbin wrote:
> The rest LGTM, but please run it through LTO review (Richard/Honza) too.

Ping?
--
Thanks, k
> 
> 	Jakub

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

[Ilya Verbin]

On 16 Oct 13:22, Jakub Jelinek wrote:
> Can you please extract the configure{,.ac}, langhooks.c and varasm.c
> bits into a separate patch?  That is preapproved for trunk right now, that
> isn't dependent on anything else.

Done.

> The rest LGTM, but please run it through LTO review (Richard/Honza) too.

Richard, is this patch OK for trunk?

Thanks,
  -- Ilya

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

[Richard Biener]

On Thu, 6 Nov 2014, Ilya Verbin wrote:

> On 16 Oct 13:22, Jakub Jelinek wrote:
> > Can you please extract the configure{,.ac}, langhooks.c and varasm.c
> > bits into a separate patch?  That is preapproved for trunk right now, that
> > isn't dependent on anything else.
> 
> Done.
> 
> > The rest LGTM, but please run it through LTO review (Richard/Honza) too.
> 
> Richard, is this patch OK for trunk?

Ok.

Thanks,
Richard.

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

[Thomas Schwinge]

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

Hi Bernd!

One question to you, as the one who has originally written this code
(<https://gcc.gnu.org/ml/gcc-patches/2014-03/msg01079.html>):

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 has been committed long ago.)

> 2014-10-02  Ilya Verbin  <ilya.verbin@intel.com>
> 	    Bernd Schmidt  <bernds@codesourcery.com>
> 	    Andrey Turetskiy  <andrey.turetskiy@intel.com>
> 	    Michael Zolotukhin  <michael.v.zolotukhin@intel.com>
> 
> gcc/
> 	* gcc.c (spec_host_machine, accel_dir_suffix): New variables.
> 	(process_command): Tweak path construction for the possibility
> 	of being configured as an offload compiler.
> 	(main): Likewise.  Look up specs in just_machine_suffix only if not
> 	ACCEL_COMPILER.  Construct OFFLOAD_TARGET_NAMES environment variable if
> 	we have OFFLOAD_TARGETS.
> 	* [...]

> --- a/gcc/gcc.c
> +++ b/gcc/gcc.c
> @@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION;
>  /* The target machine.  */
>  
>  static const char *spec_machine = DEFAULT_TARGET_MACHINE;
> +static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE;
>  
>  /* Nonzero if cross-compiling.
>     When -b is used, the value comes from the `specs' file.  */
> @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2
>     relative to the driver.  */
>  static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX;
>  
> +/* A prefix to be used when this is an accelerator compiler.  */
> +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX;
> +
>  /* Subdirectory to use for locating libraries.  Set by
>     set_multilib_dir based on the compilation options.  */
>  
> @@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count,
>      }
>  
>    gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
> -  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
> +  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
>  			    dir_separator_str, NULL);

I've noticed that internally we have a change (r447328) to »Use
spec_machine rather than spec_host_machine to build tooldir_prefix2«,
thus you reverted this last change, describing this as a »Driver paths
bugfix«, which »corrects an issue that causes an offload gcc driver to
find the wrong assembler«.

In addition to that change/reversion, one of your more recent patche
submissions, <https://gcc.gnu.org/ml/gcc-patches/2014-11/msg02655.html>,
contained the (unrelated) change to also change back to the original
spec_machine the spec_host_machine usage in the following code:

>  
>    /* Look for tools relative to the location from which the driver is
>       running, or, if that is not available, the configured prefix.  */
>    tooldir_prefix
>      = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix,
> -	      spec_machine, dir_separator_str,
> -	      spec_version, dir_separator_str, tooldir_prefix2, NULL);
> +	      spec_host_machine, dir_separator_str, spec_version,
> +	      accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL);
>    free (tooldir_prefix2);
>  
>    add_prefix (&exec_prefixes,

Which of the changes should we be making on trunk?  None at all, or just
revert the change for tooldir_prefix2 (patch variant 1), or also for
tooldir_prefix (patch variant 2), or something else?

Patch variant 1:

@@ -4266,7 +4266,7 @@ process_command (unsigned int decoded_op
     }
 
   gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
-  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
+  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
 			    dir_separator_str, NULL);
 
   /* Look for tools relative to the location from which the driver is

Patch variant 2:

@@ -4238,14 +4238,14 @@ process_command (unsigned int decoded_options_count,
     }
 
   gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
-  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
+  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
 			    dir_separator_str, NULL);
 
   /* Look for tools relative to the location from which the driver is
      running, or, if that is not available, the configured prefix.  */
   tooldir_prefix
     = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix,
-	      spec_host_machine, dir_separator_str, spec_version,
+	      spec_machine, dir_separator_str, spec_version,
 	      accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL);
   free (tooldir_prefix2);
 


Grüße,
 Thomas

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

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

[Bernd Schmidt]

On 05/12/2015 06:27 PM, Thomas Schwinge wrote:

> Patch variant 1:
>
> @@ -4266,7 +4266,7 @@ process_command (unsigned int decoded_op
>       }
>
>     gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
> -  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
> +  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
>   			    dir_separator_str, NULL);
>
>     /* Look for tools relative to the location from which the driver is

This one, definitely. The other one I can't remember. Are there any 
testsuite results for either variant?


Bernd

Help the offload gcc driver find the right assembler (was: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper)

[Thomas Schwinge]

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

Hi!

On Tue, 12 May 2015 19:05:17 +0200, Bernd Schmidt <bernds@codesourcery.com> wrote:
> > @@ -4266,7 +4266,7 @@ process_command (unsigned int decoded_op
> >       }
> >
> >     gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
> > -  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
> > +  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
> >   			    dir_separator_str, NULL);
> >
> >     /* Look for tools relative to the location from which the driver is
> 
> This one, definitely.

Indeed; now committed in r228429.  (Jim, did this missing patch cause the
thousands of FAILs you've seen?  The problem was that nvptx offloading
found the wrong assembler, and we all know what happens when you try to
process PTX assembly with an x86/PowerPC assembler...)

commit f2fd5997e0856128c0609cdc6aa4cf9867f94c41
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Fri Oct 2 21:24:50 2015 +0000

    Help the offload gcc driver find the right assembler
    
    	gcc/
    	* gcc.c (process_command): Use spec_machine rather than
    	spec_host_machine to build tooldir_prefix2.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@228429 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog |    5 +++++
 gcc/gcc.c     |    2 +-
 2 files changed, 6 insertions(+), 1 deletion(-)

diff --git gcc/ChangeLog gcc/ChangeLog
index d1235bd..7fe0196 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,8 @@
+2015-10-02  Bernd Schmidt  <bernds@codesourcery.com>
+
+	* gcc.c (process_command): Use spec_machine rather than
+	spec_host_machine to build tooldir_prefix2.
+
 2015-10-02  Thomas Schwinge  <thomas@codesourcery.com>
 	    Bernd Schmidt  <bernds@codesourcery.com>
 
diff --git gcc/gcc.c gcc/gcc.c
index db90796..7f5a36e 100644
--- gcc/gcc.c
+++ gcc/gcc.c
@@ -4472,7 +4472,7 @@ process_command (unsigned int decoded_options_count,
     }
 
   gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix));
-  tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine,
+  tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine,
 			    dir_separator_str, NULL);
 
   /* Look for tools relative to the location from which the driver is


Grüße,
 Thomas

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

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

[Thomas Schwinge]

[-- 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 --]

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

[Ilya Verbin]

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