[PATCH][RFC][Offloading] Fix PR68463

Richard Biener rguenther@suse.de
Fri Jan 15 08:15:00 GMT 2016


On Fri, 15 Jan 2016, Ilya Verbin wrote:

> Hi!
> 
> Here is my attempt to fix https://gcc.gnu.org/bugzilla/show_bug.cgi?id=68463
> 
> This patch does 2 things:
> 
> I) lto-plugin doesn't claim files which contain offload sections, but don't
> contain LTO sections.  Instead, it writes names of files with offloading to the
> temporary file and passes it to lto-wrapper as -foffload-objects=/tmp/cc...
> The order of these files in the list is very important, because ld will link
> host objects (and therefore host tables) in the following order:
>   1. Non-LTO files before the first claimed LTO file;
>   2. LTO files, after WPA-partitioning-recompilation;
>   3. Non-LTO files after the first claimed LTO file.
> To get the correct matching between host and target tables, the offload objects
> need to be reordered correspondingly before passing to the target compiler.

I think that's reasonable.

> II) The __offload_func_table, __offload_funcs_end, __offload_var_table,
> __offload_vars_end are now provided by the linker script, instead of
> crtoffload{begin,end}.o, this allows to surround all offload objects, even
> those that are not claimed by lto-plugin.
> Unfortunately it works only with ld, but doen't work with gold, because
> https://sourceware.org/bugzilla/show_bug.cgi?id=15373
> Any thoughts how to enable this linker script for gold?

The easiest way would probably to add this handling to the default
"linker script" in gold.  I don't see an easy way around requiring
changes to gold here - maybe dumping the default linker script from
bfd and injecting the rules with some scripting so you have a complete
script.  Though likely gold won't grok that result.

Really a question for Ian though.

> I used the following testcase:
> $ cat main.c
> void foo1 ();
> void foo2 ();
> void foo3 ();
> void foo4 ();
> 
> int main ()
> {
>   foo1 ();
>   foo2 ();
>   foo3 ();
>   foo4 ();
>   return 0;
> }
> 
> $ cat test.c
> #include <stdio.h>
> #include <omp.h>
> #define MAKE_FN_NAME(x) foo ## x
> #define FN_NAME(x) MAKE_FN_NAME(x)
> void FN_NAME(NUM) ()
> {
>   int x, d;
>   #pragma omp target map(from: x, d)
>     {
>       x = NUM;
>       d = omp_is_initial_device ();
>     }
>   printf ("%s:\t%s ()\tx = %d\n", d ? "HOST" : "TARGET", __FUNCTION__, x);
>   if (x != NUM)
>     printf ("--------^\n");
> }
> 
> $ gcc -DNUM=1 -c -flto test.c -o obj1.o
> $ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
> $ gcc -DNUM=3 -c test.c -o obj3.o
> $ gcc -DNUM=4 -c -flto -fopenmp test.c -o obj4.o
> $ gcc -c main.c -o main.o
> $ gcc -fopenmp obj1.o obj2.o obj3.o obj4.o main.o && ./a.out
> $ gcc -fopenmp obj2.o obj3.o obj4.o obj1.o main.o && ./a.out
> $ gcc -fopenmp obj3.o obj1.o obj2.o obj4.o main.o && ./a.out

Did you try linking an archive with both offload-but-no-LTO and
offload-and-LTO objects inside?

Thanks,
Richard.

> 
> gcc/
> 	PR driver/68463
> 	* config/i386/intelmic-mkoffload.c (generate_target_descr_file): Don't
> 	define __offload_func_table and __offload_var_table.
> 	(generate_target_offloadend_file): Remove function.
> 	(prepare_target_image): Don't call generate_target_offloadend_file.
> 	* lto-wrapper.c (offloadbegin, offloadend): Remove static vars.
> 	(offload_objects_file_name): New static var.
> 	(tool_cleanup): Remove offload_objects_file_name file.
> 	(find_offloadbeginend): Rename to ...
> 	(find_crtoffload): ... this.  Locate crtoffload.o instead of
> 	crtoffloadbegin.o and crtoffloadend.o.
> 	(run_gcc): Remove offload_argc and offload_argv.
> 	Get offload_objects_file_name from -foffload-objects=... option.
> 	Read names of object files with offload from this file, pass them to
> 	compile_images_for_offload_targets.  Call find_crtoffload instead of
> 	find_offloadbeginend.  Don't give offload files to the linker when LTO
> 	is disabled, because now they're not claimed, therefore not discarded.
> libgcc/
> 	PR driver/68463
> 	* Makefile.in (crtoffloadbegin$(objext)): Remove rule.
> 	(crtoffloadend$(objext)): Likewise.
> 	(crtoffload$(objext), link-offload-tables.x): New rules.
> 	* configure: Regenerate.
> 	* configure.ac (extra_parts): Add link-offload-tables.x if offloading is
> 	enabled, or if this is an accel compiler for intelmic.
> 	* link-offload-tables.x: New file.
> 	* offloadstuff.c: Do not define __offload_func_table,
> 	__offload_var_table, __offload_funcs_end, __offload_vars_end.
> libgomp/
> 	PR driver/68463
> 	* Makefile.in: Regenerate.
> 	* configure: Regenerate.
> 	* configure.ac (link_offload_tables): New output variable.  Set to
> 	"%Tlink-offload-tables.x" if offloading is enabled, or if this is an
> 	accel compiler for intelmic.
> 	* libgomp.spec.in (*link_gomp): Add @link_offload_tables@.
> 	* testsuite/Makefile.in: Regenerate.
> lto-plugin/
> 	PR driver/68463
> 	* lto-plugin.c (offload_files): Replace with ...
> 	(offload_files_1, offload_files_2, offload_files_3): ... this.
> 	(num_offload_files): Replace with ...
> 	(num_offload_files_1, num_offload_files_2, num_offload_files_3): ..this.
> 	(free_2): Adjust accordingly.
> 	(all_symbols_read_handler): Don't add offload files to lto_arg_ptr.
> 	Don't call free_1 for offload_files.  Write names of object files with
> 	offloading to the temporary file.  Add new option to lto_arg_ptr.
> 	(claim_file_handler): Don't claim file if it contains offload sections
> 	without LTO sections, add it to offload_files_1 or to offload_files_3.
> 	Add files with offload and LTO sections to offload_files_2.
> 
> 
> diff --git a/gcc/config/i386/intelmic-mkoffload.c b/gcc/config/i386/intelmic-mkoffload.c
> index 6a09641..82e94f1 100644
> --- a/gcc/config/i386/intelmic-mkoffload.c
> +++ b/gcc/config/i386/intelmic-mkoffload.c
> @@ -295,17 +295,12 @@ generate_target_descr_file (const char *target_compiler)
>      fatal_error (input_location, "cannot open '%s'", src_filename);
>  
>    fprintf (src_file,
> +	   "/* These symbols are provided by the linker script.  */\n"
> +	   "extern const void *const __offload_func_table[];\n"
>  	   "extern const void *const __offload_funcs_end[];\n"
> +	   "extern const void *const __offload_var_table[];\n"
>  	   "extern const void *const __offload_vars_end[];\n\n"
>  
> -	   "const void *const __offload_func_table[0]\n"
> -	   "__attribute__ ((__used__, visibility (\"hidden\"),\n"
> -	   "section (\".gnu.offload_funcs\"))) = { };\n\n"
> -
> -	   "const void *const __offload_var_table[0]\n"
> -	   "__attribute__ ((__used__, visibility (\"hidden\"),\n"
> -	   "section (\".gnu.offload_vars\"))) = { };\n\n"
> -
>  	   "const void *const __OFFLOAD_TARGET_TABLE__[]\n"
>  	   "__attribute__ ((__used__, visibility (\"hidden\"))) = {\n"
>  	   "  &__offload_func_table, &__offload_funcs_end,\n"
> @@ -342,46 +337,6 @@ generate_target_descr_file (const char *target_compiler)
>    return obj_filename;
>  }
>  
> -/* Generates object file with __offload_*_end symbols for the target
> -   library.  */
> -static const char *
> -generate_target_offloadend_file (const char *target_compiler)
> -{
> -  const char *src_filename = make_temp_file ("_target_offloadend.c");
> -  const char *obj_filename = make_temp_file ("_target_offloadend.o");
> -  temp_files[num_temps++] = src_filename;
> -  temp_files[num_temps++] = obj_filename;
> -  FILE *src_file = fopen (src_filename, "w");
> -
> -  if (!src_file)
> -    fatal_error (input_location, "cannot open '%s'", src_filename);
> -
> -  fprintf (src_file,
> -	   "const void *const __offload_funcs_end[0]\n"
> -	   "__attribute__ ((__used__, visibility (\"hidden\"),\n"
> -	   "section (\".gnu.offload_funcs\"))) = { };\n\n"
> -
> -	   "const void *const __offload_vars_end[0]\n"
> -	   "__attribute__ ((__used__, visibility (\"hidden\"),\n"
> -	   "section (\".gnu.offload_vars\"))) = { };\n");
> -  fclose (src_file);
> -
> -  struct obstack argv_obstack;
> -  obstack_init (&argv_obstack);
> -  obstack_ptr_grow (&argv_obstack, target_compiler);
> -  if (save_temps)
> -    obstack_ptr_grow (&argv_obstack, "-save-temps");
> -  if (verbose)
> -    obstack_ptr_grow (&argv_obstack, "-v");
> -  obstack_ptr_grow (&argv_obstack, "-c");
> -  obstack_ptr_grow (&argv_obstack, "-shared");
> -  obstack_ptr_grow (&argv_obstack, "-fPIC");
> -  obstack_ptr_grow (&argv_obstack, src_filename);
> -  compile_for_target (&argv_obstack, obj_filename);
> -
> -  return obj_filename;
> -}
> -
>  /* Generates object file with the host side descriptor.  */
>  static const char *
>  generate_host_descr_file (const char *host_compiler)
> @@ -469,15 +424,10 @@ prepare_target_image (const char *target_compiler, int argc, char **argv)
>  {
>    const char *target_descr_filename
>      = generate_target_descr_file (target_compiler);
> -  const char *target_offloadend_filename
> -    = generate_target_offloadend_file (target_compiler);
>  
>    char *opt1
>      = XALLOCAVEC (char, sizeof ("-Wl,") + strlen (target_descr_filename));
> -  char *opt2
> -    = XALLOCAVEC (char, sizeof ("-Wl,") + strlen (target_offloadend_filename));
>    sprintf (opt1, "-Wl,%s", target_descr_filename);
> -  sprintf (opt2, "-Wl,%s", target_offloadend_filename);
>  
>    const char *target_so_filename = make_temp_file ("_offload_intelmic.so");
>    temp_files[num_temps++] = target_so_filename;
> @@ -501,7 +451,6 @@ prepare_target_image (const char *target_compiler, int argc, char **argv)
>      }
>    if (!out_obj_filename)
>      fatal_error (input_location, "output file not specified");
> -  obstack_ptr_grow (&argv_obstack, opt2);
>    compile_for_target (&argv_obstack, target_so_filename);
>  
>    /* Run objcopy.  */
> diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c
> index bedcb79..e1d7738 100644
> --- a/gcc/lto-wrapper.c
> +++ b/gcc/lto-wrapper.c
> @@ -69,7 +69,7 @@ static char **input_names;
>  static char **output_names;
>  static char **offload_names;
>  static unsigned num_offload_targets;
> -static const char *offloadbegin, *offloadend;
> +static char *offload_objects_file_name;
>  static char *makefile;
>  
>  const char tool_name[] = "lto-wrapper";
> @@ -85,6 +85,8 @@ tool_cleanup (bool)
>      maybe_unlink (ltrans_output_file);
>    if (flto_out)
>      maybe_unlink (flto_out);
> +  if (offload_objects_file_name)
> +    maybe_unlink (offload_objects_file_name);
>    if (makefile)
>      maybe_unlink (makefile);
>    for (i = 0; i < nr; ++i)
> @@ -788,42 +790,34 @@ copy_file (const char *dest, const char *src)
>      }
>  }
>  
> -/* Find the crtoffloadbegin.o and crtoffloadend.o files in LIBRARY_PATH, make
> -   copies and store the names of the copies in offloadbegin and offloadend.  */
> +/* Find the crtoffload.o file in LIBRARY_PATH, make copy and give its name to
> +   the linker.  */
>  
>  static void
> -find_offloadbeginend (void)
> +find_crtoffload (void)
>  {
>    char **paths = NULL;
> +  const char *crtoffload;
>    const char *library_path = getenv ("LIBRARY_PATH");
>    if (!library_path)
>      return;
> -  unsigned n_paths = parse_env_var (library_path, &paths, "/crtoffloadbegin.o");
> +  unsigned n_paths = parse_env_var (library_path, &paths, "/crtoffload.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 (input_location,
> -		       "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);
> +	/* The linker will delete the filename we give it, so make a copy.  */
> +	crtoffload = make_temp_file (".crtoffload.o");
> +	copy_file (crtoffload, paths[i]);
>  	break;
>        }
>    if (i == n_paths)
> -    fatal_error (input_location,
> -		 "installation error, can't find crtoffloadbegin.o");
> +    fatal_error (input_location, "installation error, can't find crtoffload.o");
>  
>    free_array_of_ptrs ((void **) paths, n_paths);
> +
> +  printf ("%s\n", crtoffload);
>  }
>  
>  /* A subroutine of run_gcc.  Examine the open file FD for lto sections with
> @@ -918,8 +912,8 @@ run_gcc (unsigned argc, char *argv[])
>    int new_head_argc;
>    bool have_lto = false;
>    bool have_offload = false;
> -  unsigned lto_argc = 0, offload_argc = 0;
> -  char **lto_argv, **offload_argv;
> +  unsigned lto_argc = 0;
> +  char **lto_argv;
>  
>    /* Get the driver and options.  */
>    collect_gcc = getenv ("COLLECT_GCC");
> @@ -935,10 +929,9 @@ run_gcc (unsigned argc, char *argv[])
>  					&decoded_options,
>  					&decoded_options_count);
>  
> -  /* Allocate arrays for input object files with LTO or offload IL,
> +  /* Allocate array for input object files with LTO IL,
>       and for possible preceding arguments.  */
>    lto_argv = XNEWVEC (char *, argc);
> -  offload_argv = XNEWVEC (char *, argc);
>  
>    /* Look at saved options in the IL files.  */
>    for (i = 1; i < argc; ++i)
> @@ -950,6 +943,15 @@ run_gcc (unsigned argc, char *argv[])
>        int consumed;
>        char *filename = argv[i];
>  
> +      if (strncmp (argv[i], "-foffload-objects=",
> +		   sizeof ("-foffload-objects=") - 1) == 0)
> +	{
> +	  have_offload = true;
> +	  offload_objects_file_name
> +	    = argv[i] + sizeof ("-foffload-objects=") - 1;
> +	  continue;
> +	}
> +
>        if ((p = strrchr (argv[i], '@'))
>  	  && p != argv[i] 
>  	  && sscanf (p, "@%li%n", &loffset, &consumed) >= 1
> @@ -974,15 +976,6 @@ run_gcc (unsigned argc, char *argv[])
>  	  have_lto = true;
>  	  lto_argv[lto_argc++] = argv[i];
>  	}
> -
> -      if (find_and_merge_options (fd, file_offset, OFFLOAD_SECTION_NAME_PREFIX,
> -				  &offload_fdecoded_options,
> -				  &offload_fdecoded_options_count, collect_gcc))
> -	{
> -	  have_offload = true;
> -	  offload_argv[offload_argc++] = argv[i];
> -	}
> -
>        close (fd);
>      }
>  
> @@ -1081,47 +1074,83 @@ run_gcc (unsigned argc, char *argv[])
>  
>    if (have_offload)
>      {
> -      compile_images_for_offload_targets (offload_argc, offload_argv,
> +      unsigned i, num_offload_files;
> +      char **offload_argv;
> +      FILE *f;
> +
> +      f = fopen (offload_objects_file_name, "r");
> +      if (f == NULL)
> +	fatal_error (input_location, "cannot open %s: %m",
> +		     offload_objects_file_name);
> +      if (fscanf (f, "%u ", &num_offload_files) != 1)
> +	fatal_error (input_location, "cannot read %s: %m",
> +		     offload_objects_file_name);
> +      offload_argv = XNEWVEC (char *, num_offload_files);
> +
> +      /* Read names of object files with offload.  */
> +      for (i = 0; i < num_offload_files; i++)
> +	{
> +	  const unsigned piece = 32;
> +	  char *buf, *filename = XNEWVEC (char, piece);
> +	  size_t len;
> +
> +	  buf = filename;
> +cont1:
> +	  if (!fgets (buf, piece, f))
> +	    break;
> +	  len = strlen (filename);
> +	  if (filename[len - 1] != '\n')
> +	    {
> +	      filename = XRESIZEVEC (char, filename, len + piece);
> +	      buf = filename + len;
> +	      goto cont1;
> +	    }
> +	  filename[len - 1] = '\0';
> +	  offload_argv[i] = filename;
> +	}
> +      fclose (f);
> +      maybe_unlink (offload_objects_file_name);
> +      offload_objects_file_name = NULL;
> +
> +      /* Look at saved offload options in files.  */
> +      for (i = 0; i < num_offload_files; i++)
> +	{
> +	  int fd;
> +	  char *filename = offload_argv[i];
> +
> +	  fd = open (filename, O_RDONLY | O_BINARY);
> +	  if (fd == -1)
> +	    fatal_error (input_location, "cannot open %s: %m", filename);
> +	  if (!find_and_merge_options (fd, 0, OFFLOAD_SECTION_NAME_PREFIX,
> +				       &offload_fdecoded_options,
> +				       &offload_fdecoded_options_count,
> +				       collect_gcc))
> +	    fatal_error (input_location, "cannot read %s: %m", filename);
> +	  close (fd);
> +	}
> +
> +      compile_images_for_offload_targets (num_offload_files, offload_argv,
>  					  offload_fdecoded_options,
>  					  offload_fdecoded_options_count,
>  					  decoded_options,
>  					  decoded_options_count);
> +
> +      free_array_of_ptrs ((void **) offload_argv, num_offload_files);
> +
>        if (offload_names)
>  	{
> -	  find_offloadbeginend ();
> +	  find_crtoffload ();
>  	  for (i = 0; i < num_offload_targets; i++)
>  	    if (offload_names[i])
>  	      printf ("%s\n", offload_names[i]);
>  	  free_array_of_ptrs ((void **) offload_names, num_offload_targets);
>  	}
> -    }
>  
> -  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) != 0
> -	    && strncmp (argv[i], "-flinker-output=",
> -			sizeof ("-flinker-output=") - 1) != 0)
> -	  {
> -	    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 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_lto)
> +	goto finish;
>      }
>  
>    if (lto_mode == LTO_MODE_LTO)
> @@ -1351,11 +1380,7 @@ cont:
>      }
>  
>   finish:
> -  if (offloadend)
> -    printf ("%s\n", offloadend);
> -
>    XDELETE (lto_argv);
> -  XDELETE (offload_argv);
>    obstack_free (&argv_obstack, NULL);
>  }
>  
> diff --git a/libgcc/Makefile.in b/libgcc/Makefile.in
> index 570b1a7..1fdd33e 100644
> --- a/libgcc/Makefile.in
> +++ b/libgcc/Makefile.in
> @@ -994,15 +994,17 @@ crtendS$(objext): $(srcdir)/crtstuff.c
>  crtbeginT$(objext): $(srcdir)/crtstuff.c
>  	$(crt_compile) $(CRTSTUFF_T_CFLAGS) -c $< -DCRT_BEGIN -DCRTSTUFFT_O
>  
> -# crtoffloadbegin and crtoffloadend contain symbols, that mark the begin and
> +# crtoffload contains __OFFLOAD_TABLE__ symbol which points to the begin and
>  # the end of tables with addresses, required for offloading.
> -crtoffloadbegin$(objext): $(srcdir)/offloadstuff.c
> -	$(crt_compile) $(CRTSTUFF_T_CFLAGS) -c $< -DCRT_BEGIN
> -
> -crtoffloadend$(objext): $(srcdir)/offloadstuff.c
> -	$(crt_compile) $(CRTSTUFF_T_CFLAGS) -c $< -DCRT_END
> +crtoffload$(objext): $(srcdir)/offloadstuff.c
> +	$(crt_compile) $(CRTSTUFF_T_CFLAGS) -c $<
>  endif
>  
> +# This linker script provides symbols that mark the begin and the end of tables
> +# with addresses, required for offloading.
> +link-offload-tables.x: $(srcdir)/link-offload-tables.x
> +	cp $< $@
> +
>  ifeq ($(enable_vtable_verify),yes)
>  # These are used in vtable verification; see comments in source files for
>  # more details.
> diff --git a/libgcc/configure b/libgcc/configure
> index 7cf6e9b..e94ad59 100644
> --- a/libgcc/configure
> +++ b/libgcc/configure
> @@ -4829,7 +4829,14 @@ fi
>  
>  
>  if test x"$enable_offload_targets" != x; then
> -  extra_parts="${extra_parts} crtoffloadbegin.o crtoffloadend.o"
> +  extra_parts="${extra_parts} crtoffload.o link-offload-tables.x"
> +fi
> +
> +if test x"$enable_as_accelerator_for" != x; then
> +  case "${target}" in
> +    *-intelmic-* | *-intelmicemul-*)
> +      extra_parts="${extra_parts} link-offload-tables.x"
> +  esac
>  fi
>  
>  # Check if Solaris/x86 linker supports ZERO terminator unwind entries.
> diff --git a/libgcc/configure.ac b/libgcc/configure.ac
> index b96d4bc..e394b1c 100644
> --- a/libgcc/configure.ac
> +++ b/libgcc/configure.ac
> @@ -412,7 +412,14 @@ AC_SUBST(accel_dir_suffix)
>  AC_SUBST(real_host_noncanonical)
>  
>  if test x"$enable_offload_targets" != x; then
> -  extra_parts="${extra_parts} crtoffloadbegin.o crtoffloadend.o"
> +  extra_parts="${extra_parts} crtoffload.o link-offload-tables.x"
> +fi
> +
> +if test x"$enable_as_accelerator_for" != x; then
> +  case "${target}" in
> +    *-intelmic-* | *-intelmicemul-*)
> +      extra_parts="${extra_parts} link-offload-tables.x"
> +  esac
>  fi
>  
>  # Check if Solaris/x86 linker supports ZERO terminator unwind entries.
> diff --git a/libgcc/link-offload-tables.x b/libgcc/link-offload-tables.x
> new file mode 100644
> index 0000000..e7b3fb5
> --- /dev/null
> +++ b/libgcc/link-offload-tables.x
> @@ -0,0 +1,17 @@
> +SECTIONS
> +{
> +  .gnu.offload_funcs :
> +  {
> +    PROVIDE_HIDDEN (__offload_func_table = .);
> +    KEEP (*(.gnu.offload_funcs))
> +    PROVIDE_HIDDEN (__offload_funcs_end = .);
> +  }
> +
> +  .gnu.offload_vars :
> +  {
> +    PROVIDE_HIDDEN (__offload_var_table = .);
> +    KEEP (*(.gnu.offload_vars))
> +    PROVIDE_HIDDEN (__offload_vars_end = .);
> +  }
> +}
> +INSERT AFTER .data;
> diff --git a/libgcc/offloadstuff.c b/libgcc/offloadstuff.c
> index 45e89cf..eb955e3 100644
> --- a/libgcc/offloadstuff.c
> +++ b/libgcc/offloadstuff.c
> @@ -40,32 +40,13 @@ see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
>  #include "tm.h"
>  #include "libgcc_tm.h"
>  
> -#define OFFLOAD_FUNC_TABLE_SECTION_NAME ".gnu.offload_funcs"
> -#define OFFLOAD_VAR_TABLE_SECTION_NAME ".gnu.offload_vars"
> -
> -#ifdef CRT_BEGIN
> -
>  #if defined(HAVE_GAS_HIDDEN) && defined(ENABLE_OFFLOADING)
> -const void *const __offload_func_table[0]
> -  __attribute__ ((__used__, visibility ("hidden"),
> -		  section (OFFLOAD_FUNC_TABLE_SECTION_NAME))) = { };
> -const void *const __offload_var_table[0]
> -  __attribute__ ((__used__, visibility ("hidden"),
> -		  section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { };
> -#endif
> -
> -#elif defined CRT_END
> -
> -#if defined(HAVE_GAS_HIDDEN) && defined(ENABLE_OFFLOADING)
> -const void *const __offload_funcs_end[0]
> -  __attribute__ ((__used__, visibility ("hidden"),
> -		  section (OFFLOAD_FUNC_TABLE_SECTION_NAME))) = { };
> -const void *const __offload_vars_end[0]
> -  __attribute__ ((__used__, visibility ("hidden"),
> -		  section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { };
>  
> +/* These symbols are provided by the linker script.  */
>  extern const void *const __offload_func_table[];
> +extern const void *const __offload_funcs_end[];
>  extern const void *const __offload_var_table[];
> +extern const void *const __offload_vars_end[];
>  
>  const void *const __OFFLOAD_TABLE__[]
>    __attribute__ ((__visibility__ ("hidden"))) =
> @@ -73,8 +54,5 @@ const void *const __OFFLOAD_TABLE__[]
>    &__offload_func_table, &__offload_funcs_end,
>    &__offload_var_table, &__offload_vars_end
>  };
> -#endif
>  
> -#else /* ! CRT_BEGIN && ! CRT_END */
> -#error "One of CRT_BEGIN or CRT_END must be defined."
>  #endif
> diff --git a/libgomp/Makefile.in b/libgomp/Makefile.in
> index 7a1c976..dd0c861 100644
> --- a/libgomp/Makefile.in
> +++ b/libgomp/Makefile.in
> @@ -17,7 +17,7 @@
>  
>  # Plugins for offload execution, Makefile.am fragment.
>  #
> -# Copyright (C) 2014-2015 Free Software Foundation, Inc.
> +# Copyright (C) 2014-2016 Free Software Foundation, Inc.
>  #
>  # Contributed by Mentor Embedded.
>  #
> @@ -352,6 +352,7 @@ libdir = @libdir@
>  libexecdir = @libexecdir@
>  libtool_VERSION = @libtool_VERSION@
>  link_gomp = @link_gomp@
> +link_offload_tables = @link_offload_tables@
>  localedir = @localedir@
>  localstatedir = @localstatedir@
>  lt_host_flags = @lt_host_flags@
> diff --git a/libgomp/configure b/libgomp/configure
> index e2605f0..0d908ff 100755
> --- a/libgomp/configure
> +++ b/libgomp/configure
> @@ -615,6 +615,7 @@ OMP_LOCK_ALIGN
>  OMP_LOCK_SIZE
>  USE_FORTRAN_FALSE
>  USE_FORTRAN_TRUE
> +link_offload_tables
>  link_gomp
>  XLDFLAGS
>  XCFLAGS
> @@ -11121,7 +11122,7 @@ else
>    lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
>    lt_status=$lt_dlunknown
>    cat > conftest.$ac_ext <<_LT_EOF
> -#line 11124 "configure"
> +#line 11125 "configure"
>  #include "confdefs.h"
>  
>  #if HAVE_DLFCN_H
> @@ -11227,7 +11228,7 @@ else
>    lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
>    lt_status=$lt_dlunknown
>    cat > conftest.$ac_ext <<_LT_EOF
> -#line 11230 "configure"
> +#line 11231 "configure"
>  #include "confdefs.h"
>  
>  #if HAVE_DLFCN_H
> @@ -15090,7 +15091,7 @@ esac
>  
>  # Plugins for offload execution, configure.ac fragment.  -*- mode: autoconf -*-
>  #
> -# Copyright (C) 2014-2015 Free Software Foundation, Inc.
> +# Copyright (C) 2014-2016 Free Software Foundation, Inc.
>  #
>  # Contributed by Mentor Embedded.
>  #
> @@ -16478,6 +16479,20 @@ else
>  fi
>  
>  
> +# Pass link-offload-tables.x script to the linker.  It provides symbols that
> +# mark the begin and the end of tables with addresses, required for offloading.
> +link_offload_tables=
> +if test x"$enable_offload_targets" != x; then
> +  link_offload_tables="%Tlink-offload-tables.x"
> +fi
> +if test x"$enable_as_accelerator_for" != x; then
> +  case "${target}" in
> +    *-intelmic-* | *-intelmicemul-*)
> +      link_offload_tables="%Tlink-offload-tables.x"
> +  esac
> +fi
> +
> +
>   if test "$ac_cv_fc_compiler_gnu" = yes; then
>    USE_FORTRAN_TRUE=
>    USE_FORTRAN_FALSE='#'
> diff --git a/libgomp/configure.ac b/libgomp/configure.ac
> index 2e41ca8..9f8a991 100644
> --- a/libgomp/configure.ac
> +++ b/libgomp/configure.ac
> @@ -305,6 +305,20 @@ else
>  fi
>  AC_SUBST(link_gomp)
>  
> +# Pass link-offload-tables.x script to the linker.  It provides symbols that
> +# mark the begin and the end of tables with addresses, required for offloading.
> +link_offload_tables=
> +if test x"$enable_offload_targets" != x; then
> +  link_offload_tables="%Tlink-offload-tables.x"
> +fi
> +if test x"$enable_as_accelerator_for" != x; then
> +  case "${target}" in
> +    *-intelmic-* | *-intelmicemul-*)
> +      link_offload_tables="%Tlink-offload-tables.x"
> +  esac
> +fi
> +AC_SUBST(link_offload_tables)
> +
>  AM_CONDITIONAL([USE_FORTRAN], [test "$ac_cv_fc_compiler_gnu" = yes])
>  
>  # ??? 2006-01-24: Paulo committed to asking autoconf folk to document
> diff --git a/libgomp/libgomp.spec.in b/libgomp/libgomp.spec.in
> index 5651603..6a946c4 100644
> --- a/libgomp/libgomp.spec.in
> +++ b/libgomp/libgomp.spec.in
> @@ -1,3 +1,3 @@
>  # This spec file is read by gcc when linking.  It is used to specify the
>  # standard libraries we need in order to link with libgomp.
> -*link_gomp: @link_gomp@
> +*link_gomp: @link_gomp@ @link_offload_tables@
> diff --git a/libgomp/testsuite/Makefile.in b/libgomp/testsuite/Makefile.in
> index c25d21f..a3982bf 100644
> --- a/libgomp/testsuite/Makefile.in
> +++ b/libgomp/testsuite/Makefile.in
> @@ -208,6 +208,7 @@ libdir = @libdir@
>  libexecdir = @libexecdir@
>  libtool_VERSION = @libtool_VERSION@
>  link_gomp = @link_gomp@
> +link_offload_tables = @link_offload_tables@
>  localedir = @localedir@
>  localstatedir = @localstatedir@
>  lt_host_flags = @lt_host_flags@
> diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c
> index 0a6a767..a62c31e 100644
> --- a/lto-plugin/lto-plugin.c
> +++ b/lto-plugin/lto-plugin.c
> @@ -152,8 +152,14 @@ static ld_plugin_add_symbols add_symbols;
>  static struct plugin_file_info *claimed_files = NULL;
>  static unsigned int num_claimed_files = 0;
>  
> -static struct plugin_file_info *offload_files = NULL;
> -static unsigned int num_offload_files = 0;
> +/* Lists of files with offloading.  We need 3 of them to maintain the correct
> +   order, otherwise host and target tables with addresses wouldn't match.  */
> +static char **offload_files_1;
> +static char **offload_files_2;
> +static char **offload_files_3;
> +static unsigned num_offload_files_1;
> +static unsigned num_offload_files_2;
> +static unsigned num_offload_files_3;
>  
>  static char **output_files = NULL;
>  static unsigned int num_output_files = 0;
> @@ -351,14 +357,6 @@ free_2 (void)
>        free (info->name);
>      }
>  
> -  for (i = 0; i < num_offload_files; i++)
> -    {
> -      struct plugin_file_info *info = &offload_files[i];
> -      struct plugin_symtab *symtab = &info->symtab;
> -      free (symtab->aux);
> -      free (info->name);
> -    }
> -
>    for (i = 0; i < num_output_files; i++)
>      free (output_files[i]);
>    free (output_files);
> @@ -367,9 +365,17 @@ free_2 (void)
>    claimed_files = NULL;
>    num_claimed_files = 0;
>  
> -  free (offload_files);
> -  offload_files = NULL;
> -  num_offload_files = 0;
> +  for (i = 0; i < num_offload_files_1; i++)
> +    free (offload_files_1[i]);
> +  for (i = 0; i < num_offload_files_2; i++)
> +    free (offload_files_2[i]);
> +  for (i = 0; i < num_offload_files_3; i++)
> +    free (offload_files_3[i]);
> +  free (offload_files_1);
> +  free (offload_files_2);
> +  free (offload_files_3);
> +  offload_files_1 = offload_files_2 = offload_files_3 = NULL;
> +  num_offload_files_1 = num_offload_files_2 = num_offload_files_3 = 0;
>  
>    free (arguments_file_name);
>    arguments_file_name = NULL;
> @@ -625,11 +631,12 @@ static enum ld_plugin_status
>  all_symbols_read_handler (void)
>  {
>    unsigned i;
> -  unsigned num_lto_args
> -    = num_claimed_files + num_offload_files + lto_wrapper_num_args + 2;
> +  unsigned num_lto_args = num_claimed_files + lto_wrapper_num_args + 3;
>    char **lto_argv;
>    const char *linker_output_str;
>    const char **lto_arg_ptr;
> +  unsigned num_offload_files
> +    = num_offload_files_1 + num_offload_files_2 + num_offload_files_3;
>    if (num_claimed_files + num_offload_files == 0)
>      return LDPS_OK;
>  
> @@ -646,7 +653,6 @@ all_symbols_read_handler (void)
>    write_resolution ();
>  
>    free_1 (claimed_files, num_claimed_files);
> -  free_1 (offload_files, num_offload_files);
>  
>    for (i = 0; i < lto_wrapper_num_args; i++)
>      *lto_arg_ptr++ = lto_wrapper_argv[i];
> @@ -671,16 +677,40 @@ all_symbols_read_handler (void)
>        break;
>      }
>    *lto_arg_ptr++ = xstrdup (linker_output_str);
> -  for (i = 0; i < num_claimed_files; i++)
> -    {
> -      struct plugin_file_info *info = &claimed_files[i];
>  
> -      *lto_arg_ptr++ = info->name;
> +  if (num_offload_files > 0)
> +    {
> +      FILE *f;
> +      char *arg;
> +      char *offload_objects_file_name;
> +
> +      offload_objects_file_name = make_temp_file ("");
> +      check (offload_objects_file_name, LDPL_FATAL,
> +	     "Failed to generate a temporary file name");
> +      f = fopen (offload_objects_file_name, "w");
> +      check (f, LDPL_FATAL, "could not open file with offload objects");
> +      fprintf (f, "%u\n", num_offload_files);
> +
> +      /* Names of files with offloading are written in the following order:
> +	 1. Non-LTO files before the first claimed LTO file;
> +	 2. LTO files;
> +	 3. Non-LTO files after the first claimed LTO file.  */
> +      for (i = 0; i < num_offload_files_1; i++)
> +	fprintf (f, "%s\n", offload_files_1[i]);
> +      for (i = 0; i < num_offload_files_2; i++)
> +	fprintf (f, "%s\n", offload_files_2[i]);
> +      for (i = 0; i < num_offload_files_3; i++)
> +	fprintf (f, "%s\n", offload_files_3[i]);
> +      fclose (f);
> +
> +      arg = concat ("-foffload-objects=", offload_objects_file_name, NULL);
> +      check (arg, LDPL_FATAL, "could not allocate");
> +      *lto_arg_ptr++ = arg;
>      }
>  
> -  for (i = 0; i < num_offload_files; i++)
> +  for (i = 0; i < num_claimed_files; i++)
>      {
> -      struct plugin_file_info *info = &offload_files[i];
> +      struct plugin_file_info *info = &claimed_files[i];
>  
>        *lto_arg_ptr++ = info->name;
>      }
> @@ -1007,18 +1037,37 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed)
>  	xrealloc (claimed_files,
>  		  num_claimed_files * sizeof (struct plugin_file_info));
>        claimed_files[num_claimed_files - 1] = lto_file;
> +
> +      *claimed = 1;
>      }
>  
> -  if (obj.found == 0 && obj.offload == 1)
> +  if (obj.offload == 1)
>      {
> -      num_offload_files++;
> -      offload_files =
> -	xrealloc (offload_files,
> -		  num_offload_files * sizeof (struct plugin_file_info));
> -      offload_files[num_offload_files - 1] = lto_file;
> -    }
> +      char ***arr;
> +      unsigned *num;
> +      if (num_claimed_files == 0)
> +	{
> +	  /* Offload Non-LTO file before the first claimed LTO file.  */
> +	  arr = &offload_files_1;
> +	  num = &num_offload_files_1;
> +	}
> +      else if (*claimed)
> +	{
> +	  /* Offload LTO file.  */
> +	  arr = &offload_files_2;
> +	  num = &num_offload_files_2;
> +	}
> +      else
> +	{
> +	  /* Offload Non-LTO file after the first claimed LTO file.  */
> +	  arr = &offload_files_3;
> +	  num = &num_offload_files_3;
> +	}
>  
> -  *claimed = 1;
> +      (*num)++;
> +      *arr = xrealloc (*arr, *num * sizeof (char *));
> +      (*arr)[*num - 1] = xstrdup (lto_file.name);
> +    }
>  
>    goto cleanup;
> 
> 
> Thanks,
>   -- Ilya
> 
> 

-- 
Richard Biener <rguenther@suse.de>
SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB 21284 (AG Nuernberg)



More information about the Gcc-patches mailing list