This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [PATCH][RFC][Offloading] Fix PR68463
- From: Richard Biener <rguenther at suse dot de>
- To: Ilya Verbin <iverbin at gmail dot com>
- Cc: Jakub Jelinek <jakub at redhat dot com>, bschmidt at redhat dot com, gcc-patches at gcc dot gnu dot org, kirill dot yukhin at gmail dot com, thomas at codesourcery dot com, iant at google dot com
- Date: Fri, 15 Jan 2016 09:15:01 +0100 (CET)
- Subject: Re: [PATCH][RFC][Offloading] Fix PR68463
- Authentication-results: sourceware.org; auth=none
- References: <20160114212601 dot GA48907 at msticlxl57 dot ims dot intel dot com>
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)