+2016-02-25 Ilya Verbin <ilya.verbin@intel.com>
+
+ 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 <hubicka@ucw.cz>
PR ipa/69630
#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
%{" 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
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
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";
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)
}
}
-/* 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);
}
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");
&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)
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
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);
}
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)
{
}
finish:
- if (offloadend)
- printf ("%s\n", offloadend);
-
XDELETE (lto_argv);
- XDELETE (offload_argv);
obstack_free (&argv_obstack, NULL);
}
+2016-02-25 Ilya Verbin <ilya.verbin@intel.com>
+
+ 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 <jcmvbkbc@gmail.com>
* config/xtensa/ieee754-df.S (__muldf3_aux, __divdf3_aux): Add
$(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)
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.
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.
#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))) = { };
__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"))) =
&__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
+2016-02-25 Ilya Verbin <ilya.verbin@intel.com>
+
+ PR driver/68463
+ * testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c: Remove.
+
2016-02-23 Thomas Schwinge <thomas@codesourcery.com>
* oacc-parallel.c (GOACC_parallel_keyed): Initialize dims.
+++ /dev/null
-/* { 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;
-}
-
+2016-02-25 Ilya Verbin <ilya.verbin@intel.com>
+
+ 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 <mliska@suse.cz>
* lto-plugin.c (all_symbols_read_handler): Assign default
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. */
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;
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);
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);
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;
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];
break;
}
*lto_arg_ptr++ = xstrdup (linker_output_str);
- for (i = 0; i < num_claimed_files; i++)
+
+ if (num_offload_files > 0)
{
- struct plugin_file_info *info = &claimed_files[i];
+ 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);
- *lto_arg_ptr++ = info->name;
+ 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;
}
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;
projects / gcc.git / commitdiff