diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 9daf0c01ce7..21d304d1cc0 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,23 @@ +2016-02-25 Ilya Verbin + + PR driver/68463 + * config/gnu-user.h (CRTOFFLOADBEGIN): Define. Add crtoffloadbegin.o if + offloading is enabled and -fopenacc or -fopenmp is specified. + (CRTOFFLOADEND): Likewise. + (GNU_USER_TARGET_STARTFILE_SPEC): Add CRTOFFLOADBEGIN. + (GNU_USER_TARGET_ENDFILE_SPEC): Add CRTOFFLOADEND. + * 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): Replace with ... + (find_crtoffloadtable): ... this. + (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. Don't call find_offloadbeginend and + don't pass offloadbegin and offloadend to the linker. Don't pass + offload non-LTO files to the linker, because now they're not claimed. + 2016-02-25 Jan Hubicka PR ipa/69630 diff --git a/gcc/config/gnu-user.h b/gcc/config/gnu-user.h index 2f1bbccb055..b0bf40a954f 100644 --- a/gcc/config/gnu-user.h +++ b/gcc/config/gnu-user.h @@ -35,6 +35,14 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see #undef ASM_APP_OFF #define ASM_APP_OFF "#NO_APP\n" +#if ENABLE_OFFLOADING == 1 +#define CRTOFFLOADBEGIN "%{fopenacc|fopenmp:crtoffloadbegin%O%s}" +#define CRTOFFLOADEND "%{fopenacc|fopenmp:crtoffloadend%O%s}" +#else +#define CRTOFFLOADBEGIN "" +#define CRTOFFLOADEND "" +#endif + /* Provide a STARTFILE_SPEC appropriate for GNU userspace. Here we add the GNU userspace magical crtbegin.o file (see crtstuff.c) which provides part of the support for getting C++ file-scope static @@ -49,14 +57,16 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see %{" NO_PIE_SPEC ":crtbegin.o%s}} \ %{fvtable-verify=none:%s; \ fvtable-verify=preinit:vtv_start_preinit.o%s; \ - fvtable-verify=std:vtv_start.o%s}" + fvtable-verify=std:vtv_start.o%s} \ + " CRTOFFLOADBEGIN #else #define GNU_USER_TARGET_STARTFILE_SPEC \ "%{!shared: %{pg|p|profile:gcrt1.o%s;:crt1.o%s}} \ crti.o%s %{static:crtbeginT.o%s;shared|pie:crtbeginS.o%s;:crtbegin.o%s} \ %{fvtable-verify=none:%s; \ fvtable-verify=preinit:vtv_start_preinit.o%s; \ - fvtable-verify=std:vtv_start.o%s}" + fvtable-verify=std:vtv_start.o%s} \ + " CRTOFFLOADBEGIN #endif #undef STARTFILE_SPEC #define STARTFILE_SPEC GNU_USER_TARGET_STARTFILE_SPEC @@ -73,13 +83,15 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see fvtable-verify=preinit:vtv_end_preinit.o%s; \ fvtable-verify=std:vtv_end.o%s} \ %{shared:crtendS.o%s;: %{" PIE_SPEC ":crtendS.o%s} \ - %{" NO_PIE_SPEC ":crtend.o%s}} crtn.o%s" + %{" NO_PIE_SPEC ":crtend.o%s}} crtn.o%s \ + " CRTOFFLOADEND #else #define GNU_USER_TARGET_ENDFILE_SPEC \ "%{fvtable-verify=none:%s; \ fvtable-verify=preinit:vtv_end_preinit.o%s; \ fvtable-verify=std:vtv_end.o%s} \ - %{shared|pie:crtendS.o%s;:crtend.o%s} crtn.o%s" + %{shared|pie:crtendS.o%s;:crtend.o%s} crtn.o%s \ + " CRTOFFLOADEND #endif #undef ENDFILE_SPEC #define ENDFILE_SPEC GNU_USER_TARGET_ENDFILE_SPEC diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c index ed20b4ee95e..f240812f952 100644 --- a/gcc/lto-wrapper.c +++ b/gcc/lto-wrapper.c @@ -68,7 +68,7 @@ static unsigned int nr; static char **input_names; static char **output_names; static char **offload_names; -static const char *offloadbegin, *offloadend; +static char *offload_objects_file_name; static char *makefile; const char tool_name[] = "lto-wrapper"; @@ -84,6 +84,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) @@ -840,40 +842,32 @@ 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 crtoffloadtable.o file in LIBRARY_PATH, make copy and pass name of + the copy to the linker. */ static void -find_offloadbeginend (void) +find_crtoffloadtable (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 n_paths = parse_env_var (library_path, &paths, "/crtoffloadtable.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. */ + char *crtoffloadtable = make_temp_file (".crtoffloadtable.o"); + copy_file (crtoffloadtable, paths[i]); + printf ("%s\n", crtoffloadtable); + XDELETEVEC (crtoffloadtable); break; } if (i == n_paths) fatal_error (input_location, - "installation error, can't find crtoffloadbegin.o"); + "installation error, can't find crtoffloadtable.o"); free_array_of_ptrs ((void **) paths, n_paths); } @@ -970,8 +964,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"); @@ -987,10 +981,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) @@ -1002,6 +995,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 @@ -1026,15 +1028,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); } @@ -1133,47 +1126,102 @@ 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 = XCNEWVEC (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); + if (offload_argv[num_offload_files - 1] == NULL) + fatal_error (input_location, "invalid format of %s", + offload_objects_file_name); + 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++) + { + char *p; + long loffset; + int fd, consumed; + off_t file_offset = 0; + char *filename = offload_argv[i]; + + if ((p = strrchr (offload_argv[i], '@')) + && p != offload_argv[i] + && sscanf (p, "@%li%n", &loffset, &consumed) >= 1 + && strlen (p) == (unsigned int) consumed) + { + filename = XNEWVEC (char, p - offload_argv[i] + 1); + memcpy (filename, offload_argv[i], p - offload_argv[i]); + filename[p - offload_argv[i]] = '\0'; + file_offset = (off_t) loffset; + } + 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, file_offset, + OFFLOAD_SECTION_NAME_PREFIX, + &offload_fdecoded_options, + &offload_fdecoded_options_count, + collect_gcc)) + fatal_error (input_location, "cannot read %s: %m", filename); + close (fd); + if (filename != offload_argv[i]) + XDELETEVEC (filename); + } + + 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_crtoffloadtable (); 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) != 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; - } + goto finish; if (lto_mode == LTO_MODE_LTO) { @@ -1402,11 +1450,7 @@ cont: } finish: - if (offloadend) - printf ("%s\n", offloadend); - XDELETE (lto_argv); - XDELETE (offload_argv); obstack_free (&argv_obstack, NULL); } diff --git a/libgcc/ChangeLog b/libgcc/ChangeLog index 4020e235bf1..36d10b0450f 100644 --- a/libgcc/ChangeLog +++ b/libgcc/ChangeLog @@ -1,3 +1,13 @@ +2016-02-25 Ilya Verbin + + PR driver/68463 + * Makefile.in (crtoffloadtable$(objext)): New rule. + * configure.ac (extra_parts): Add crtoffloadtable$(objext) if + enable_offload_targets is not empty. + * configure: Regenerate. + * offloadstuff.c: Move __OFFLOAD_TABLE__ from crtoffloadend to + crtoffloadtable. + 2016-02-17 Max Filippov * config/xtensa/ieee754-df.S (__muldf3_aux, __divdf3_aux): Add diff --git a/libgcc/Makefile.in b/libgcc/Makefile.in index 570b1a7da07..f09b39b0e85 100644 --- a/libgcc/Makefile.in +++ b/libgcc/Makefile.in @@ -995,12 +995,16 @@ 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 -# the end of tables with addresses, required for offloading. +# the end of tables with addresses, required for offloading. crtoffloadtable +# contains the array with addresses of those symbols. 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 + +crtoffloadtable$(objext): $(srcdir)/offloadstuff.c + $(crt_compile) $(CRTSTUFF_T_CFLAGS) -c $< -DCRT_TABLE endif ifeq ($(enable_vtable_verify),yes) diff --git a/libgcc/configure b/libgcc/configure index de8c13c4ba3..f3f360512c4 100644 --- a/libgcc/configure +++ b/libgcc/configure @@ -4835,7 +4835,7 @@ fi if test x"$enable_offload_targets" != x; then - extra_parts="${extra_parts} crtoffloadbegin.o crtoffloadend.o" + extra_parts="${extra_parts} crtoffloadbegin.o crtoffloadend.o crtoffloadtable.o" fi # Check if Solaris/x86 linker supports ZERO terminator unwind entries. diff --git a/libgcc/configure.ac b/libgcc/configure.ac index 860a5f58100..897259e62bc 100644 --- a/libgcc/configure.ac +++ b/libgcc/configure.ac @@ -418,7 +418,7 @@ 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} crtoffloadbegin.o crtoffloadend.o crtoffloadtable.o" fi # Check if Solaris/x86 linker supports ZERO terminator unwind entries. diff --git a/libgcc/offloadstuff.c b/libgcc/offloadstuff.c index 45e89cf9df1..4ab639721cb 100644 --- a/libgcc/offloadstuff.c +++ b/libgcc/offloadstuff.c @@ -40,23 +40,22 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see #include "tm.h" #include "libgcc_tm.h" +#if defined(HAVE_GAS_HIDDEN) && ENABLE_OFFLOADING == 1 + #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))) = { }; @@ -64,8 +63,12 @@ const void *const __offload_vars_end[0] __attribute__ ((__used__, visibility ("hidden"), section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { }; +#elif defined CRT_TABLE + extern const void *const __offload_func_table[]; extern const void *const __offload_var_table[]; +extern const void *const __offload_funcs_end[]; +extern const void *const __offload_vars_end[]; const void *const __OFFLOAD_TABLE__[] __attribute__ ((__visibility__ ("hidden"))) = @@ -73,8 +76,9 @@ const void *const __OFFLOAD_TABLE__[] &__offload_func_table, &__offload_funcs_end, &__offload_var_table, &__offload_vars_end }; + +#else /* ! CRT_BEGIN && ! CRT_END && ! CRT_TABLE */ +#error "One of CRT_BEGIN, CRT_END or CRT_TABLE must be defined." #endif -#else /* ! CRT_BEGIN && ! CRT_END */ -#error "One of CRT_BEGIN or CRT_END must be defined." #endif diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index e6a708255b0..3cf28601f0d 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2016-02-25 Ilya Verbin + + PR driver/68463 + * testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c: Remove. + 2016-02-23 Thomas Schwinge * oacc-parallel.c (GOACC_parallel_keyed): Initialize dims. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c deleted file mode 100644 index eea8c7e4fa2..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c +++ /dev/null @@ -1,19 +0,0 @@ -/* { dg-do run { target { openacc_nvidia_accel_selected && lto } } } */ -/* { dg-additional-options "-flto -fno-use-linker-plugin" } */ - -/* Worker and vector size checks. Picked an outrageously large - value. */ - -int main () -{ -#pragma acc parallel num_workers (2<<20) /* { dg-error "using num_workers" } */ - { - } - -#pragma acc parallel vector_length (2<<20) /* { dg-error "using vector_length" } */ - { - } - - return 0; -} - diff --git a/lto-plugin/ChangeLog b/lto-plugin/ChangeLog index 4ebccb71df8..d2cd98675bc 100644 --- a/lto-plugin/ChangeLog +++ b/lto-plugin/ChangeLog @@ -1,3 +1,17 @@ +2016-02-25 Ilya Verbin + + PR driver/68463 + * lto-plugin.c (struct plugin_offload_file): New. + (offload_files): Change type. + (offload_files_last, offload_files_last_obj): New. + (offload_files_last_lto): New. + (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. If it contains offload sections, add to the list. + 2016-01-15 Martin Liska * lto-plugin.c (all_symbols_read_handler): Assign default diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c index 1ed0f0863c0..51afc528d05 100644 --- a/lto-plugin/lto-plugin.c +++ b/lto-plugin/lto-plugin.c @@ -129,6 +129,14 @@ struct plugin_file_info struct plugin_symtab conflicts; }; +/* List item with name of the file with offloading. */ + +struct plugin_offload_file +{ + char *name; + struct plugin_offload_file *next; +}; + /* Until ASM_OUTPUT_LABELREF can be hookized and decoupled from stdio file streams, we do simple label translation here. */ @@ -152,8 +160,16 @@ 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; +/* List of files with offloading. */ +static struct plugin_offload_file *offload_files; +/* Last file in the list. */ +static struct plugin_offload_file *offload_files_last; +/* Last non-archive file in the list. */ +static struct plugin_offload_file *offload_files_last_obj; +/* Last LTO file in the list. */ +static struct plugin_offload_file *offload_files_last_lto; +/* Total number of files with offloading. */ +static unsigned num_offload_files; static char **output_files = NULL; static unsigned int num_output_files = 0; @@ -351,14 +367,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,8 +375,12 @@ free_2 (void) claimed_files = NULL; num_claimed_files = 0; - free (offload_files); - offload_files = NULL; + while (offload_files) + { + struct plugin_offload_file *ofld = offload_files; + offload_files = offload_files->next; + free (ofld); + } num_offload_files = 0; free (arguments_file_name); @@ -625,8 +637,7 @@ 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 = NULL; const char **lto_arg_ptr; @@ -646,7 +657,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,6 +681,35 @@ all_symbols_read_handler (void) break; } *lto_arg_ptr++ = xstrdup (linker_output_str); + + if (num_offload_files > 0) + { + FILE *f; + char *arg; + char *offload_objects_file_name; + struct plugin_offload_file *ofld; + + offload_objects_file_name = make_temp_file (".ofldlist"); + 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); + + /* Skip the dummy item at the start of the list. */ + ofld = offload_files->next; + while (ofld) + { + fprintf (f, "%s\n", ofld->name); + ofld = ofld->next; + } + 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_claimed_files; i++) { struct plugin_file_info *info = &claimed_files[i]; @@ -678,13 +717,6 @@ all_symbols_read_handler (void) *lto_arg_ptr++ = info->name; } - for (i = 0; i < num_offload_files; i++) - { - struct plugin_file_info *info = &offload_files[i]; - - *lto_arg_ptr++ = info->name; - } - *lto_arg_ptr++ = NULL; exec_lto_wrapper (lto_argv); @@ -1007,18 +1039,72 @@ 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 (offload_files == NULL) { - num_offload_files++; - offload_files = - xrealloc (offload_files, - num_offload_files * sizeof (struct plugin_file_info)); - offload_files[num_offload_files - 1] = lto_file; + /* Add dummy item to the start of the list. */ + offload_files = xmalloc (sizeof (struct plugin_offload_file)); + offload_files->name = NULL; + offload_files->next = NULL; + offload_files_last = offload_files; } - *claimed = 1; + /* If this is an LTO file without offload, and it is the first LTO file, save + the pointer to the last offload file in the list. Further offload LTO + files will be inserted after it, if any. */ + if (*claimed && obj.offload == 0 && offload_files_last_lto == NULL) + offload_files_last_lto = offload_files_last; + + if (obj.offload == 1) + { + /* Add file to the list. The order must be exactly the same as the final + order after recompilation and linking, otherwise host and target tables + with addresses wouldn't match. If a static library contains both LTO + and non-LTO objects, ld and gold link them in a different order. */ + struct plugin_offload_file *ofld + = xmalloc (sizeof (struct plugin_offload_file)); + ofld->name = lto_file.name; + ofld->next = NULL; + + if (*claimed && offload_files_last_lto == NULL && file->offset != 0 + && gold_version == -1) + { + /* ld only: insert first LTO file from the archive after the last real + object file immediately preceding the archive, or at the begin of + the list if there was no real objects before archives. */ + if (offload_files_last_obj != NULL) + { + ofld->next = offload_files_last_obj->next; + offload_files_last_obj->next = ofld; + } + else + { + ofld->next = offload_files->next; + offload_files->next = ofld; + } + } + else if (*claimed && offload_files_last_lto != NULL) + { + /* Insert LTO file after the last LTO file in the list. */ + ofld->next = offload_files_last_lto->next; + offload_files_last_lto->next = ofld; + } + else + /* Add non-LTO file or first non-archive LTO file to the end of the + list. */ + offload_files_last->next = ofld; + + if (ofld->next == NULL) + offload_files_last = ofld; + if (file->offset == 0) + offload_files_last_obj = ofld; + if (*claimed) + offload_files_last_lto = ofld; + num_offload_files++; + } goto cleanup;