* [PATCH][RFC][Offloading] Fix PR68463
@ 2016-01-14 21:26 Ilya Verbin
2016-01-15 8:15 ` Richard Biener
0 siblings, 1 reply; 23+ messages in thread
From: Ilya Verbin @ 2016-01-14 21:26 UTC (permalink / raw)
To: rguenther, jakub, bschmidt, gcc-patches; +Cc: kirill.yukhin, thomas
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.
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?
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
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
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-01-14 21:26 [PATCH][RFC][Offloading] Fix PR68463 Ilya Verbin
@ 2016-01-15 8:15 ` Richard Biener
2016-01-18 20:34 ` Ilya Verbin
0 siblings, 1 reply; 23+ messages in thread
From: Richard Biener @ 2016-01-15 8:15 UTC (permalink / raw)
To: Ilya Verbin
Cc: Jakub Jelinek, bschmidt, gcc-patches, kirill.yukhin, thomas, iant
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)
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-01-15 8:15 ` Richard Biener
@ 2016-01-18 20:34 ` Ilya Verbin
2016-01-19 8:57 ` Richard Biener
0 siblings, 1 reply; 23+ messages in thread
From: Ilya Verbin @ 2016-01-18 20:34 UTC (permalink / raw)
To: Richard Biener, Jakub Jelinek
Cc: bschmidt, gcc-patches, kirill.yukhin, thomas, iant
On Fri, Jan 15, 2016 at 09:15:01 +0100, Richard Biener wrote:
> On Fri, 15 Jan 2016, Ilya Verbin wrote:
> > 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.
Or the gcc driver can add crtoffload{begin,end}.o, but the problem is that it
can't determine whether the program contains offloading or not. So it can add
them to all -fopenmp/-fopenacc programs, if the compiler was configured with
--enable-offload-targets=... The overhead would be about 340 bytes for
binaries which doesn't use offloading. Is this acceptable? (Jakub?)
> > 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?
No. And it didn't work, because archives are handled by ld a bit differently.
I will fix it. Thanks! From ld/ldlang.c:
/* Find the insert point for the plugin's replacement files. We
place them after the first claimed real object file, or if the
first claimed object is an archive member, after the last real
object file immediately preceding the archive.
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-01-18 20:34 ` Ilya Verbin
@ 2016-01-19 8:57 ` Richard Biener
2016-01-19 9:36 ` Jakub Jelinek
0 siblings, 1 reply; 23+ messages in thread
From: Richard Biener @ 2016-01-19 8:57 UTC (permalink / raw)
To: Ilya Verbin
Cc: Jakub Jelinek, bschmidt, gcc-patches, kirill.yukhin, thomas, iant
On Mon, 18 Jan 2016, Ilya Verbin wrote:
> On Fri, Jan 15, 2016 at 09:15:01 +0100, Richard Biener wrote:
> > On Fri, 15 Jan 2016, Ilya Verbin wrote:
> > > 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.
>
> Or the gcc driver can add crtoffload{begin,end}.o, but the problem is that it
> can't determine whether the program contains offloading or not. So it can add
> them to all -fopenmp/-fopenacc programs, if the compiler was configured with
> --enable-offload-targets=... The overhead would be about 340 bytes for
> binaries which doesn't use offloading. Is this acceptable? (Jakub?)
Can lto-wrapper add them as plugin outputs? Or does that wreck ordering?
Richard.
>
> > > 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?
>
> No. And it didn't work, because archives are handled by ld a bit differently.
> I will fix it. Thanks! From ld/ldlang.c:
>
> /* Find the insert point for the plugin's replacement files. We
> place them after the first claimed real object file, or if the
> first claimed object is an archive member, after the last real
> object file immediately preceding the archive.
>
> -- Ilya
>
>
--
Richard Biener <rguenther@suse.de>
SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB 21284 (AG Nuernberg)
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-01-19 8:57 ` Richard Biener
@ 2016-01-19 9:36 ` Jakub Jelinek
2016-01-19 13:32 ` Ilya Verbin
0 siblings, 1 reply; 23+ messages in thread
From: Jakub Jelinek @ 2016-01-19 9:36 UTC (permalink / raw)
To: Richard Biener
Cc: Ilya Verbin, bschmidt, gcc-patches, kirill.yukhin, thomas, iant
On Tue, Jan 19, 2016 at 09:57:01AM +0100, Richard Biener wrote:
> On Mon, 18 Jan 2016, Ilya Verbin wrote:
>
> > On Fri, Jan 15, 2016 at 09:15:01 +0100, Richard Biener wrote:
> > > On Fri, 15 Jan 2016, Ilya Verbin wrote:
> > > > 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.
> >
> > Or the gcc driver can add crtoffload{begin,end}.o, but the problem is that it
> > can't determine whether the program contains offloading or not. So it can add
> > them to all -fopenmp/-fopenacc programs, if the compiler was configured with
> > --enable-offload-targets=... The overhead would be about 340 bytes for
> > binaries which doesn't use offloading. Is this acceptable? (Jakub?)
>
> Can lto-wrapper add them as plugin outputs? Or does that wreck ordering?
Yeah, if that would work, it would be certainly appreciated, one thing is
wasting .text space and relocations in all -fopenmp programs (for -fopenacc
programs one kind of assumes there will be some offloading in there),
another one some extra constructor/destructor or what that would be even
worse.
Jakub
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-01-19 9:36 ` Jakub Jelinek
@ 2016-01-19 13:32 ` Ilya Verbin
2016-02-10 17:20 ` Ilya Verbin
0 siblings, 1 reply; 23+ messages in thread
From: Ilya Verbin @ 2016-01-19 13:32 UTC (permalink / raw)
To: Jakub Jelinek, Richard Biener
Cc: bschmidt, gcc-patches, kirill.yukhin, thomas, iant
On Tue, Jan 19, 2016 at 10:36:28 +0100, Jakub Jelinek wrote:
> On Tue, Jan 19, 2016 at 09:57:01AM +0100, Richard Biener wrote:
> > On Mon, 18 Jan 2016, Ilya Verbin wrote:
> > > On Fri, Jan 15, 2016 at 09:15:01 +0100, Richard Biener wrote:
> > > > On Fri, 15 Jan 2016, Ilya Verbin wrote:
> > > > > 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.
> > >
> > > Or the gcc driver can add crtoffload{begin,end}.o, but the problem is that it
> > > can't determine whether the program contains offloading or not. So it can add
> > > them to all -fopenmp/-fopenacc programs, if the compiler was configured with
> > > --enable-offload-targets=... The overhead would be about 340 bytes for
> > > binaries which doesn't use offloading. Is this acceptable? (Jakub?)
> >
> > Can lto-wrapper add them as plugin outputs? Or does that wreck ordering?
Currently it's implemented this way, but it will not work after my patch,
because e.g. offload-without-lto.o and offload-with-lto.o will be linked in
this order:
offload-without-lto.o, crtoffloadbegin.o, offload-with-lto.o, crtoffloadend.o
^^^^^^^^^^^^^^^^^^^^^
(will be not claimed by the plugin)
But we need this one:
crtoffloadbegin.o, offload-without-lto.o, offload-with-lto.o, crtoffloadend.o
> Yeah, if that would work, it would be certainly appreciated, one thing is
> wasting .text space and relocations in all -fopenmp programs (for -fopenacc
> programs one kind of assumes there will be some offloading in there),
> another one some extra constructor/destructor or what that would be even
> worse.
They contain only 5 symbols, without constructors/destructors.
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-01-19 13:32 ` Ilya Verbin
@ 2016-02-10 17:20 ` Ilya Verbin
2016-02-19 14:53 ` Jakub Jelinek
0 siblings, 1 reply; 23+ messages in thread
From: Ilya Verbin @ 2016-02-10 17:20 UTC (permalink / raw)
To: Jakub Jelinek
Cc: bschmidt, gcc-patches, kirill.yukhin, thomas, Richard Biener
Hi!
On Tue, Jan 19, 2016 at 16:32:13 +0300, Ilya Verbin wrote:
> On Tue, Jan 19, 2016 at 10:36:28 +0100, Jakub Jelinek wrote:
> > On Tue, Jan 19, 2016 at 09:57:01AM +0100, Richard Biener wrote:
> > > On Mon, 18 Jan 2016, Ilya Verbin wrote:
> > > > On Fri, Jan 15, 2016 at 09:15:01 +0100, Richard Biener wrote:
> > > > > On Fri, 15 Jan 2016, Ilya Verbin wrote:
> > > > > > 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.
> > > >
> > > > Or the gcc driver can add crtoffload{begin,end}.o, but the problem is that it
> > > > can't determine whether the program contains offloading or not. So it can add
> > > > them to all -fopenmp/-fopenacc programs, if the compiler was configured with
> > > > --enable-offload-targets=... The overhead would be about 340 bytes for
> > > > binaries which doesn't use offloading. Is this acceptable? (Jakub?)
> > >
> > > Can lto-wrapper add them as plugin outputs? Or does that wreck ordering?
>
> Currently it's implemented this way, but it will not work after my patch,
> because e.g. offload-without-lto.o and offload-with-lto.o will be linked in
> this order:
> offload-without-lto.o, crtoffloadbegin.o, offload-with-lto.o, crtoffloadend.o
> ^^^^^^^^^^^^^^^^^^^^^
> (will be not claimed by the plugin)
>
> But we need this one:
> crtoffloadbegin.o, offload-without-lto.o, offload-with-lto.o, crtoffloadend.o
>
> > Yeah, if that would work, it would be certainly appreciated, one thing is
> > wasting .text space and relocations in all -fopenmp programs (for -fopenacc
> > programs one kind of assumes there will be some offloading in there),
> > another one some extra constructor/destructor or what that would be even
> > worse.
>
> They contain only 5 symbols, without constructors/destructors.
This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
I couldn't think of a better solution...
Tested using the testcase from the previous mail, e.g.:
$ gcc -DNUM=1 -c -fopenmp test.c -o obj1.o
$ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
$ gcc -DNUM=3 -c -fopenmp test.c -o obj3.o
$ gcc -DNUM=4 -c -fopenmp test.c -o obj4.o -flto
$ gcc -DNUM=5 -c -fopenmp test.c -o obj5.o
$ gcc -DNUM=6 -c -fopenmp test.c -o obj6.o -flto
$ gcc -DNUM=7 -c -fopenmp test.c -o obj7.o
$ gcc-ar -cvq libtest.a obj3.o obj4.o obj5.o
$ gcc -fopenmp main.c obj1.o obj2.o libtest.a obj6.o obj7.o
And other combinations.
gcc/
PR driver/68463
* config/gnu-user.h (GNU_USER_TARGET_STARTFILE_SPEC): Add
crtoffloadbegin.o for -fopenacc/-fopenmp if it exists.
(GNU_USER_TARGET_ENDFILE_SPEC): Add crtoffloadend.o for
-fopenacc/-fopenmp if it exists.
* lto-wrapper.c (offloadbegin, offloadend): Remove static vars.
(offload_objects_file_name): New static var.
(tool_cleanup): Remove offload_objects_file_name file.
(copy_file): Remove function.
(find_offloadbeginend): Remove function.
(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.
lto-plugin/
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.
diff --git a/gcc/config/gnu-user.h b/gcc/config/gnu-user.h
index 2f1bbcc..2fdb63c 100644
--- a/gcc/config/gnu-user.h
+++ b/gcc/config/gnu-user.h
@@ -49,14 +49,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} \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
#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} \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
#endif
#undef STARTFILE_SPEC
#define STARTFILE_SPEC GNU_USER_TARGET_STARTFILE_SPEC
@@ -73,13 +75,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 \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
#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 \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
#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 ed20b4e..f2914d0 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)
@@ -818,66 +820,6 @@ compile_images_for_offload_targets (unsigned in_argc, char *in_argv[],
free_array_of_ptrs ((void **) names, num_targets);
}
-/* Copy a file from SRC to DEST. */
-
-static void
-copy_file (const char *dest, const char *src)
-{
- FILE *d = fopen (dest, "wb");
- FILE *s = fopen (src, "rb");
- char buffer[512];
- while (!feof (s))
- {
- size_t len = fread (buffer, 1, 512, s);
- if (ferror (s) != 0)
- fatal_error (input_location, "reading input file");
- if (len > 0)
- {
- fwrite (buffer, 1, len, d);
- if (ferror (d) != 0)
- fatal_error (input_location, "writing output file");
- }
- }
-}
-
-/* Find the crtoffloadbegin.o and crtoffloadend.o files in LIBRARY_PATH, make
- copies and store the names of the copies in offloadbegin and offloadend. */
-
-static void
-find_offloadbeginend (void)
-{
- char **paths = NULL;
- const char *library_path = getenv ("LIBRARY_PATH");
- if (!library_path)
- return;
- unsigned n_paths = parse_env_var (library_path, &paths, "/crtoffloadbegin.o");
-
- unsigned i;
- for (i = 0; i < n_paths; i++)
- if (access_check (paths[i], R_OK) == 0)
- {
- size_t len = strlen (paths[i]);
- char *tmp = xstrdup (paths[i]);
- strcpy (paths[i] + len - strlen ("begin.o"), "end.o");
- if (access_check (paths[i], R_OK) != 0)
- fatal_error (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);
- break;
- }
- if (i == n_paths)
- fatal_error (input_location,
- "installation error, can't find crtoffloadbegin.o");
-
- free_array_of_ptrs ((void **) paths, n_paths);
-}
-
/* A subroutine of run_gcc. Examine the open file FD for lto sections with
name prefix PREFIX, at FILE_OFFSET, and store any options we find in OPTS
and OPT_COUNT. Return true if we found a matchingn section, false
@@ -970,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");
@@ -987,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)
@@ -1002,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
@@ -1026,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);
}
@@ -1133,47 +1074,101 @@ 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 ();
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 +1397,7 @@ cont:
}
finish:
- if (offloadend)
- printf ("%s\n", offloadend);
-
XDELETE (lto_argv);
- XDELETE (offload_argv);
obstack_free (&argv_obstack, NULL);
}
diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c
index 1ed0f08..9aba151 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,16 +681,37 @@ all_symbols_read_handler (void)
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);
+
+ ofld = offload_files;
+ 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;
}
@@ -1007,19 +1038,63 @@ 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)
{
+ struct plugin_offload_file *ofld
+ = xmalloc (sizeof (struct plugin_offload_file));
+ ofld->name = lto_file.name;
+ ofld->next = NULL;
+
+ if (offload_files == NULL)
+ offload_files = ofld;
+
+ /* 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. */
+ 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 if (offload_files != ofld)
+ {
+ ofld->next = offload_files;
+ offload_files = 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 if (offload_files_last != NULL)
+ {
+ /* 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++;
- offload_files =
- xrealloc (offload_files,
- num_offload_files * sizeof (struct plugin_file_info));
- offload_files[num_offload_files - 1] = lto_file;
}
- *claimed = 1;
-
goto cleanup;
err:
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-10 17:20 ` Ilya Verbin
@ 2016-02-19 14:53 ` Jakub Jelinek
2016-02-19 17:58 ` Mike Stump
2016-02-20 10:55 ` Ilya Verbin
0 siblings, 2 replies; 23+ messages in thread
From: Jakub Jelinek @ 2016-02-19 14:53 UTC (permalink / raw)
To: Ilya Verbin; +Cc: bschmidt, gcc-patches, kirill.yukhin, thomas, Richard Biener
On Wed, Feb 10, 2016 at 08:19:34PM +0300, Ilya Verbin wrote:
> This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
> I couldn't think of a better solution...
> Tested using the testcase from the previous mail, e.g.:
>
> $ gcc -DNUM=1 -c -fopenmp test.c -o obj1.o
> $ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
> $ gcc -DNUM=3 -c -fopenmp test.c -o obj3.o
> $ gcc -DNUM=4 -c -fopenmp test.c -o obj4.o -flto
> $ gcc -DNUM=5 -c -fopenmp test.c -o obj5.o
> $ gcc -DNUM=6 -c -fopenmp test.c -o obj6.o -flto
> $ gcc -DNUM=7 -c -fopenmp test.c -o obj7.o
> $ gcc-ar -cvq libtest.a obj3.o obj4.o obj5.o
> $ gcc -fopenmp main.c obj1.o obj2.o libtest.a obj6.o obj7.o
>
> And other combinations.
Looking at this, I think I have no problem with crtoffloadbegin.o being
included in all -fopenmp/-fopenacc linked programs/shared libraries,
that just defines the symbols and nothing else.
I have no problem with the
__offload_funcs_end/__offload_vars_end part of crtoffloadend.o being
included too.
But, I really don't like __OFFLOAD_TABLE__ being added to all programs, that
wastes real space in data (rodata or relro?) section, and dynamic
relocations.
So, perhaps, can we split offloadstuff.c into 3 objects instead of 2,
crtoffload{begin,end,table}.o*, where the last one would be what
defines __OFFLOAD_TABLE__, and add the last one only by the linker
plugin/lto-wrapper/whatever, if any input objects had any offloading stuff
in it?
Jakub
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-19 14:53 ` Jakub Jelinek
@ 2016-02-19 17:58 ` Mike Stump
2016-02-20 10:55 ` Ilya Verbin
1 sibling, 0 replies; 23+ messages in thread
From: Mike Stump @ 2016-02-19 17:58 UTC (permalink / raw)
To: Jakub Jelinek
Cc: Ilya Verbin, bschmidt, gcc-patches, kirill.yukhin, thomas,
Richard Biener
On Feb 19, 2016, at 6:53 AM, Jakub Jelinek <jakub@redhat.com> wrote:
> Looking at this, I think I have no problem with crtoffloadbegin.o being
> included in all -fopenmp/-fopenacc linked programs/shared libraries,
:-) I have a problem with just the normal init path in most executables. It adds a ton of stuff that can be empty at the bottom. I sometimes wonder if we boosted it to -flto, and then let lto see the size of the table, and put all the init code under an early if (count) { do the init stuff; }, then given the count, lto can then just remove it all, reliably.
If the openmp people want to experiment with -flto and see if they can make the whole thing disappear that way, it might be worth considering.
But, yes, I agree, hard to want yet more included by default that just won’t go away.
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-19 14:53 ` Jakub Jelinek
2016-02-19 17:58 ` Mike Stump
@ 2016-02-20 10:55 ` Ilya Verbin
2016-02-22 10:59 ` Jakub Jelinek
2016-02-22 15:13 ` Thomas Schwinge
1 sibling, 2 replies; 23+ messages in thread
From: Ilya Verbin @ 2016-02-20 10:55 UTC (permalink / raw)
To: Jakub Jelinek, Thomas Schwinge, Richard Biener
Cc: bschmidt, gcc-patches, kirill.yukhin
On Fri, Feb 19, 2016 at 15:53:08 +0100, Jakub Jelinek wrote:
> On Wed, Feb 10, 2016 at 08:19:34PM +0300, Ilya Verbin wrote:
> > This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
> > I couldn't think of a better solution...
> > Tested using the testcase from the previous mail, e.g.:
> >
> > $ gcc -DNUM=1 -c -fopenmp test.c -o obj1.o
> > $ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
> > $ gcc -DNUM=3 -c -fopenmp test.c -o obj3.o
> > $ gcc -DNUM=4 -c -fopenmp test.c -o obj4.o -flto
> > $ gcc -DNUM=5 -c -fopenmp test.c -o obj5.o
> > $ gcc -DNUM=6 -c -fopenmp test.c -o obj6.o -flto
> > $ gcc -DNUM=7 -c -fopenmp test.c -o obj7.o
> > $ gcc-ar -cvq libtest.a obj3.o obj4.o obj5.o
> > $ gcc -fopenmp main.c obj1.o obj2.o libtest.a obj6.o obj7.o
> >
> > And other combinations.
>
> Looking at this, I think I have no problem with crtoffloadbegin.o being
> included in all -fopenmp/-fopenacc linked programs/shared libraries,
> that just defines the symbols and nothing else.
> I have no problem with the
> __offload_funcs_end/__offload_vars_end part of crtoffloadend.o being
> included too.
> But, I really don't like __OFFLOAD_TABLE__ being added to all programs, that
> wastes real space in data (rodata or relro?) section, and dynamic
> relocations.
> So, perhaps, can we split offloadstuff.c into 3 objects instead of 2,
> crtoffload{begin,end,table}.o*, where the last one would be what
> defines __OFFLOAD_TABLE__, and add the last one only by the linker
> plugin/lto-wrapper/whatever, if any input objects had any offloading stuff
> in it?
Done. Bootstrapped and regtested, lto-bootstrap in progress.
Thomas, could you please test it using nvptx, including the testcase with static
libraries?
Could this patch be considered for stage4? On the one hand, this is not a
regression. On the other hand, it fixes quite serious issues, and it shouldn't
affect non-offloading configurations.
gcc/
PR driver/68463
* config/gnu-user.h (GNU_USER_TARGET_STARTFILE_SPEC): Add
crtoffloadbegin.o for -fopenacc/-fopenmp if it exists.
(GNU_USER_TARGET_ENDFILE_SPEC): Add crtoffloadend.o for
-fopenacc/-fopenmp if it exists.
* 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.
libgcc/
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.
lto-plugin/
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.
diff --git a/gcc/config/gnu-user.h b/gcc/config/gnu-user.h
index 2f1bbcc..2fdb63c 100644
--- a/gcc/config/gnu-user.h
+++ b/gcc/config/gnu-user.h
@@ -49,14 +49,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} \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
#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} \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
#endif
#undef STARTFILE_SPEC
#define STARTFILE_SPEC GNU_USER_TARGET_STARTFILE_SPEC
@@ -73,13 +75,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 \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
#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 \
+ %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
#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 ed20b4e..f240812 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/Makefile.in b/libgcc/Makefile.in
index 570b1a7..f09b39b 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 de8c13c..f3f3605 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 860a5f5..897259e 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 45e89cf..a4ea3ac 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) && defined(ENABLE_OFFLOADING)
+
#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/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c
index 1ed0f08..35cb63a 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,16 +681,37 @@ all_symbols_read_handler (void)
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);
+
+ 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;
}
@@ -1007,18 +1038,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;
Thanks,
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-20 10:55 ` Ilya Verbin
@ 2016-02-22 10:59 ` Jakub Jelinek
2016-02-22 15:13 ` Thomas Schwinge
1 sibling, 0 replies; 23+ messages in thread
From: Jakub Jelinek @ 2016-02-22 10:59 UTC (permalink / raw)
To: Ilya Verbin
Cc: Thomas Schwinge, Richard Biener, bschmidt, gcc-patches, kirill.yukhin
On Sat, Feb 20, 2016 at 01:54:20PM +0300, Ilya Verbin wrote:
> gcc/
> PR driver/68463
> * config/gnu-user.h (GNU_USER_TARGET_STARTFILE_SPEC): Add
> crtoffloadbegin.o for -fopenacc/-fopenmp if it exists.
> (GNU_USER_TARGET_ENDFILE_SPEC): Add crtoffloadend.o for
> -fopenacc/-fopenmp if it exists.
> * 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.
> libgcc/
> 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.
> lto-plugin/
> 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.
LGTM.
Jakub
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-20 10:55 ` Ilya Verbin
2016-02-22 10:59 ` Jakub Jelinek
@ 2016-02-22 15:13 ` Thomas Schwinge
2016-02-22 18:08 ` Ilya Verbin
` (2 more replies)
1 sibling, 3 replies; 23+ messages in thread
From: Thomas Schwinge @ 2016-02-22 15:13 UTC (permalink / raw)
To: Ilya Verbin, Jakub Jelinek
Cc: bschmidt, gcc-patches, kirill.yukhin, Richard Biener, Tom de Vries
Hi!
On Sat, 20 Feb 2016 13:54:20 +0300, Ilya Verbin <iverbin@gmail.com> wrote:
> On Fri, Feb 19, 2016 at 15:53:08 +0100, Jakub Jelinek wrote:
> > On Wed, Feb 10, 2016 at 08:19:34PM +0300, Ilya Verbin wrote:
> > > This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
> > > I couldn't think of a better solution...
> > > Tested using the testcase from the previous mail, e.g.:
> > >
> > > $ gcc -DNUM=1 -c -fopenmp test.c -o obj1.o
> > > $ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
> > > $ gcc -DNUM=3 -c -fopenmp test.c -o obj3.o
> > > $ gcc -DNUM=4 -c -fopenmp test.c -o obj4.o -flto
> > > $ gcc -DNUM=5 -c -fopenmp test.c -o obj5.o
> > > $ gcc -DNUM=6 -c -fopenmp test.c -o obj6.o -flto
> > > $ gcc -DNUM=7 -c -fopenmp test.c -o obj7.o
> > > $ gcc-ar -cvq libtest.a obj3.o obj4.o obj5.o
> > > $ gcc -fopenmp main.c obj1.o obj2.o libtest.a obj6.o obj7.o
> > >
> > > And other combinations.
> Thomas, could you please test it using nvptx
It mostly ;-) works. With nvptx offloading enabled (which you don't
have, do you?), I'm seeing one test case regress:
[-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 9)
[-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 13)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
(Same for C++.) That testcase, just recently added by Tom in r233237
"Handle -fdiagnostics-color in lto", specifies 'dg-additional-options
"-flto -fno-use-linker-plugin"'. Is that now an unsupported
combination/configuration? (I have not yet looked in detail, but it
appears as if the offloading compilers are no longer being run for
-fno-use-linker-plugin.)
> including the testcase with static
> libraries?
Works in my manual testing if I work around the following issue:
> --- a/gcc/config/gnu-user.h
> +++ b/gcc/config/gnu-user.h
> @@ -49,14 +49,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} \
> + %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
(..., and similar for others.) The if-exists spec function only works
for absolute paths (I have not researched, why?), so it won't locate the
files for relative -Bbuild-gcc/[...] prefixes, and linking will fail:
/tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x0): undefined reference to `__offload_func_table'
/tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x8): undefined reference to `__offload_funcs_end'
/tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x10): undefined reference to `__offload_var_table'
/tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x18): undefined reference to `__offload_vars_end'
If I use the absolute -B$PWD/build-gcc/[...], it works. (But there is no
requirement for -B prefixes to be absolute, as far as I know.) Why not
make it a hard error, though, if these files are missing? Can we use
something like (untested pseudo-patch):
+#ifdef ENABLE_OFFLOADING
+# define CRTOFFLOADBEGIN "%{fopenacc|fopenmp:%:crtoffloadbegin%O%s}"
+#else
+# define CRTOFFLOADBEGIN ""
+#endif
@@ -49,14 +49,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 ")}"
I have not verified your patch's logic in detail (arcane...) ;-) so just
two drive-by comments:
> #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} \
> + %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
> #endif
> #undef STARTFILE_SPEC
> #define STARTFILE_SPEC GNU_USER_TARGET_STARTFILE_SPEC
> @@ -73,13 +75,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 \
> + %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
> #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 \
> + %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
> #endif
> #undef ENDFILE_SPEC
> #define ENDFILE_SPEC GNU_USER_TARGET_ENDFILE_SPEC
I guess we currently don't have to care about offloading configurations
not using the gnu-user.h file in which you modified the
STARTFILE_SPEC/ENDFILE_SPEC?
> --- a/lto-plugin/lto-plugin.c
> +++ b/lto-plugin/lto-plugin.c
> @@ -671,16 +681,37 @@ all_symbols_read_handler (void)
> + if (num_offload_files > 0)
> {
> + [...]
> + struct plugin_offload_file *ofld;
> + [...]
> + ofld = offload_files->next;
> + while (ofld)
> + {
> + fprintf (f, "%s\n", ofld->name);
> + ofld = ofld->next;
> + }
To the casual reader, skipping the first offload_files looks like a
off-by-one error, so I suggest you add a comment "Skip the dummy item at
the start of the list.", or similar.
Grüße
Thomas
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-22 15:13 ` Thomas Schwinge
@ 2016-02-22 18:08 ` Ilya Verbin
2016-02-23 7:37 ` Tom de Vries
2016-02-24 19:30 ` Ilya Verbin
2016-05-10 15:41 ` [PATCH] Apply fix for PR68463 to RS6000 James Norris
2 siblings, 1 reply; 23+ messages in thread
From: Ilya Verbin @ 2016-02-22 18:08 UTC (permalink / raw)
To: Thomas Schwinge
Cc: Jakub Jelinek, Bernd Schmidt, gcc, Kirill Yukhin, Richard Biener,
Tom de Vries
2016-02-22 18:13 GMT+03:00 Thomas Schwinge <thomas@codesourcery.com>:
> On Sat, 20 Feb 2016 13:54:20 +0300, Ilya Verbin <iverbin@gmail.com> wrote:
>> On Fri, Feb 19, 2016 at 15:53:08 +0100, Jakub Jelinek wrote:
>> > On Wed, Feb 10, 2016 at 08:19:34PM +0300, Ilya Verbin wrote:
>> > > This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
>> > > I couldn't think of a better solution...
>> > > Tested using the testcase from the previous mail, e.g.:
>> > >
>> > > $ gcc -DNUM=1 -c -fopenmp test.c -o obj1.o
>> > > $ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
>> > > $ gcc -DNUM=3 -c -fopenmp test.c -o obj3.o
>> > > $ gcc -DNUM=4 -c -fopenmp test.c -o obj4.o -flto
>> > > $ gcc -DNUM=5 -c -fopenmp test.c -o obj5.o
>> > > $ gcc -DNUM=6 -c -fopenmp test.c -o obj6.o -flto
>> > > $ gcc -DNUM=7 -c -fopenmp test.c -o obj7.o
>> > > $ gcc-ar -cvq libtest.a obj3.o obj4.o obj5.o
>> > > $ gcc -fopenmp main.c obj1.o obj2.o libtest.a obj6.o obj7.o
>> > >
>> > > And other combinations.
>
>> Thomas, could you please test it using nvptx
>
> It mostly ;-) works. With nvptx offloading enabled (which you don't
> have, do you?), I'm seeing one test case regress:
>
> [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 9)
> [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 13)
> PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>
> (Same for C++.) That testcase, just recently added by Tom in r233237
> "Handle -fdiagnostics-color in lto", specifies 'dg-additional-options
> "-flto -fno-use-linker-plugin"'. Is that now an unsupported
> combination/configuration? (I have not yet looked in detail, but it
> appears as if the offloading compilers are no longer being run for
> -fno-use-linker-plugin.)
Yes, it's really hard to fix the "lto + non-lto objects" issue for
no-use-linker-plugin LTO path. In this patch lto-plugin prepares a
list of objects files with offloading and passes it to lto-wrapper, so
I believe we should consider offloading without lto-plugin as
unsupported. I'll update wiki when the patch will be committed.
>> including the testcase with static
>> libraries?
>
> Works in my manual testing if I work around the following issue:
>
>> --- a/gcc/config/gnu-user.h
>> +++ b/gcc/config/gnu-user.h
>> @@ -49,14 +49,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} \
>> + %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
>
> (..., and similar for others.) The if-exists spec function only works
> for absolute paths (I have not researched, why?), so it won't locate the
> files for relative -Bbuild-gcc/[...] prefixes, and linking will fail:
>
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x0): undefined reference to `__offload_func_table'
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x8): undefined reference to `__offload_funcs_end'
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x10): undefined reference to `__offload_var_table'
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x18): undefined reference to `__offload_vars_end'
>
> If I use the absolute -B$PWD/build-gcc/[...], it works. (But there is no
> requirement for -B prefixes to be absolute, as far as I know.) Why not
> make it a hard error, though, if these files are missing? Can we use
> something like (untested pseudo-patch):
>
> +#ifdef ENABLE_OFFLOADING
> +# define CRTOFFLOADBEGIN "%{fopenacc|fopenmp:%:crtoffloadbegin%O%s}"
> +#else
> +# define CRTOFFLOADBEGIN ""
> +#endif
>
> @@ -49,14 +49,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 ")}"
OK, I'll replace if-exists with ifdef ENABLE_OFFLOADING.
> I have not verified your patch's logic in detail (arcane...) ;-) so just
> two drive-by comments:
>
>> #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} \
>> + %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
>> #endif
>> #undef STARTFILE_SPEC
>> #define STARTFILE_SPEC GNU_USER_TARGET_STARTFILE_SPEC
>> @@ -73,13 +75,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 \
>> + %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
>> #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 \
>> + %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
>> #endif
>> #undef ENDFILE_SPEC
>> #define ENDFILE_SPEC GNU_USER_TARGET_ENDFILE_SPEC
>
> I guess we currently don't have to care about offloading configurations
> not using the gnu-user.h file in which you modified the
> STARTFILE_SPEC/ENDFILE_SPEC?
I think so.
>> --- a/lto-plugin/lto-plugin.c
>> +++ b/lto-plugin/lto-plugin.c
>
>> @@ -671,16 +681,37 @@ all_symbols_read_handler (void)
>
>> + if (num_offload_files > 0)
>> {
>> + [...]
>> + struct plugin_offload_file *ofld;
>> + [...]
>> + ofld = offload_files->next;
>> + while (ofld)
>> + {
>> + fprintf (f, "%s\n", ofld->name);
>> + ofld = ofld->next;
>> + }
>
> To the casual reader, skipping the first offload_files looks like a
> off-by-one error, so I suggest you add a comment "Skip the dummy item at
> the start of the list.", or similar.
OK.
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-22 18:08 ` Ilya Verbin
@ 2016-02-23 7:37 ` Tom de Vries
2016-02-24 16:13 ` Thomas Schwinge
0 siblings, 1 reply; 23+ messages in thread
From: Tom de Vries @ 2016-02-23 7:37 UTC (permalink / raw)
To: Ilya Verbin, Thomas Schwinge
Cc: Jakub Jelinek, Bernd Schmidt, gcc, Kirill Yukhin, Richard Biener
On 22/02/16 19:07, Ilya Verbin wrote:
> 2016-02-22 18:13 GMT+03:00 Thomas Schwinge<thomas@codesourcery.com>:
>> >On Sat, 20 Feb 2016 13:54:20 +0300, Ilya Verbin<iverbin@gmail.com> wrote:
>>> >>On Fri, Feb 19, 2016 at 15:53:08 +0100, Jakub Jelinek wrote:
>>>> >> >On Wed, Feb 10, 2016 at 08:19:34PM +0300, Ilya Verbin wrote:
>>>>> >> > >This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
>>>>> >> > >I couldn't think of a better solution...
>>>>> >> > >Tested using the testcase from the previous mail, e.g.:
>>>>> >> > >
>>>>> >> > >$ gcc -DNUM=1 -c -fopenmp test.c -o obj1.o
>>>>> >> > >$ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
>>>>> >> > >$ gcc -DNUM=3 -c -fopenmp test.c -o obj3.o
>>>>> >> > >$ gcc -DNUM=4 -c -fopenmp test.c -o obj4.o -flto
>>>>> >> > >$ gcc -DNUM=5 -c -fopenmp test.c -o obj5.o
>>>>> >> > >$ gcc -DNUM=6 -c -fopenmp test.c -o obj6.o -flto
>>>>> >> > >$ gcc -DNUM=7 -c -fopenmp test.c -o obj7.o
>>>>> >> > >$ gcc-ar -cvq libtest.a obj3.o obj4.o obj5.o
>>>>> >> > >$ gcc -fopenmp main.c obj1.o obj2.o libtest.a obj6.o obj7.o
>>>>> >> > >
>>>>> >> > >And other combinations.
>> >
>>> >>Thomas, could you please test it using nvptx
>> >
>> >It mostly;-) works. With nvptx offloading enabled (which you don't
>> >have, do you?), I'm seeing one test case regress:
>> >
>> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 9)
>> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 13)
>> > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
>> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>> >
>> >(Same for C++.) That testcase, just recently added by Tom in r233237
>> >"Handle -fdiagnostics-color in lto", specifies 'dg-additional-options
>> >"-flto -fno-use-linker-plugin"'. Is that now an unsupported
>> >combination/configuration? (I have not yet looked in detail, but it
>> >appears as if the offloading compilers are no longer being run for
>> >-fno-use-linker-plugin.)
> Yes, it's really hard to fix the "lto + non-lto objects" issue for
> no-use-linker-plugin LTO path. In this patch lto-plugin prepares a
> list of objects files with offloading and passes it to lto-wrapper, so
> I believe we should consider offloading without lto-plugin as
> unsupported. I'll update wiki when the patch will be committed.
>
Shouldn't we error (or at least warn) then if we compile a file
containing an offload construct with fopenacc/fopenmp and
-fno-use-linker-plugin?
Thanks,
- Tom
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-23 7:37 ` Tom de Vries
@ 2016-02-24 16:13 ` Thomas Schwinge
2016-02-24 16:30 ` Ilya Verbin
0 siblings, 1 reply; 23+ messages in thread
From: Thomas Schwinge @ 2016-02-24 16:13 UTC (permalink / raw)
To: Tom de Vries, Ilya Verbin
Cc: Jakub Jelinek, Bernd Schmidt, gcc, Kirill Yukhin, Richard Biener
Hi!
On Tue, 23 Feb 2016 08:37:07 +0100, Tom de Vries <Tom_deVries@mentor.com> wrote:
> On 22/02/16 19:07, Ilya Verbin wrote:
> > 2016-02-22 18:13 GMT+03:00 Thomas Schwinge<thomas@codesourcery.com>:
> >> >On Sat, 20 Feb 2016 13:54:20 +0300, Ilya Verbin<iverbin@gmail.com> wrote:
> >>> >>On Fri, Feb 19, 2016 at 15:53:08 +0100, Jakub Jelinek wrote:
> >>>> >> >On Wed, Feb 10, 2016 at 08:19:34PM +0300, Ilya Verbin wrote:
> >>>>> >> > >This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
> >>> >>Thomas, could you please test it using nvptx
> >> >
> >> >It mostly;-) works. With nvptx offloading enabled (which you don't
> >> >have, do you?), I'm seeing one test case regress:
> >> >
> >> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 9)
> >> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 13)
> >> > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
> >> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
> >> >
> >> >(Same for C++.) That testcase, just recently added by Tom in r233237
> >> >"Handle -fdiagnostics-color in lto", specifies 'dg-additional-options
> >> >"-flto -fno-use-linker-plugin"'. Is that now an unsupported
> >> >combination/configuration? (I have not yet looked in detail, but it
> >> >appears as if the offloading compilers are no longer being run for
> >> >-fno-use-linker-plugin.)
> > Yes, it's really hard to fix the "lto + non-lto objects" issue for
> > no-use-linker-plugin LTO path. In this patch lto-plugin prepares a
> > list of objects files with offloading and passes it to lto-wrapper, so
> > I believe we should consider offloading without lto-plugin as
> > unsupported. I'll update wiki when the patch will be committed.
Aha, I see. I guess there's no point in keeping offloading supported for
the -fno-lto (default) with -fno-use-linker-plugin configuration?
Ilya, then please remove
libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c as part of
your patch, unless Tom thinks it should be changed to a -flto test, but
without -fno-use-linker-plugin?
> Shouldn't we error (or at least warn) then if we compile a file
> containing an offload construct with fopenacc/fopenmp and
> -fno-use-linker-plugin?
Yes, that makes sense to me, too. (Note that, as I understand it,
-fno-use-linker-plugin may also be the default for certain GCC
configurations...) Aside from spec stuff in gcc/gcc.c relating to
LINK_PLUGIN_SPEC, I see there's some code in
gcc/gcc.c:driver::maybe_run_linker evaluating the three possible values
of HAVE_LTO_PLUGIN, but I have not yet thought about how and where to
conditionalize the diagnostic if attempting to do offloading in an
unsupported (-fno-use-linker-plugin) configuration.
Grüße
Thomas
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-24 16:13 ` Thomas Schwinge
@ 2016-02-24 16:30 ` Ilya Verbin
0 siblings, 0 replies; 23+ messages in thread
From: Ilya Verbin @ 2016-02-24 16:30 UTC (permalink / raw)
To: Thomas Schwinge
Cc: Tom de Vries, Jakub Jelinek, Bernd Schmidt, gcc, Kirill Yukhin,
Richard Biener
On Wed, Feb 24, 2016 at 17:13:35 +0100, Thomas Schwinge wrote:
> On Tue, 23 Feb 2016 08:37:07 +0100, Tom de Vries <Tom_deVries@mentor.com> wrote:
> > On 22/02/16 19:07, Ilya Verbin wrote:
> > > 2016-02-22 18:13 GMT+03:00 Thomas Schwinge<thomas@codesourcery.com>:
> > >> >On Sat, 20 Feb 2016 13:54:20 +0300, Ilya Verbin<iverbin@gmail.com> wrote:
> > >>> >>On Fri, Feb 19, 2016 at 15:53:08 +0100, Jakub Jelinek wrote:
> > >>>> >> >On Wed, Feb 10, 2016 at 08:19:34PM +0300, Ilya Verbin wrote:
> > >>>>> >> > >This patch adds crtoffload{begin,end}.o to all -fopenmp programs, if they exist.
>
> > >>> >>Thomas, could you please test it using nvptx
> > >> >
> > >> >It mostly;-) works. With nvptx offloading enabled (which you don't
> > >> >have, do you?), I'm seeing one test case regress:
> > >> >
> > >> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 9)
> > >> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for errors, line 13)
> > >> > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
> > >> > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
> > >> >
> > >> >(Same for C++.) That testcase, just recently added by Tom in r233237
> > >> >"Handle -fdiagnostics-color in lto", specifies 'dg-additional-options
> > >> >"-flto -fno-use-linker-plugin"'. Is that now an unsupported
> > >> >combination/configuration? (I have not yet looked in detail, but it
> > >> >appears as if the offloading compilers are no longer being run for
> > >> >-fno-use-linker-plugin.)
> > > Yes, it's really hard to fix the "lto + non-lto objects" issue for
> > > no-use-linker-plugin LTO path. In this patch lto-plugin prepares a
> > > list of objects files with offloading and passes it to lto-wrapper, so
> > > I believe we should consider offloading without lto-plugin as
> > > unsupported. I'll update wiki when the patch will be committed.
>
> Aha, I see. I guess there's no point in keeping offloading supported for
> the -fno-lto (default) with -fno-use-linker-plugin configuration?
>
> Ilya, then please remove
> libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c as part of
> your patch, unless Tom thinks it should be changed to a -flto test, but
> without -fno-use-linker-plugin?
OK.
> > Shouldn't we error (or at least warn) then if we compile a file
> > containing an offload construct with fopenacc/fopenmp and
> > -fno-use-linker-plugin?
>
> Yes, that makes sense to me, too. (Note that, as I understand it,
> -fno-use-linker-plugin may also be the default for certain GCC
> configurations...) Aside from spec stuff in gcc/gcc.c relating to
> LINK_PLUGIN_SPEC, I see there's some code in
> gcc/gcc.c:driver::maybe_run_linker evaluating the three possible values
> of HAVE_LTO_PLUGIN, but I have not yet thought about how and where to
> conditionalize the diagnostic if attempting to do offloading in an
> unsupported (-fno-use-linker-plugin) configuration.
To print this error someone has to detect that at least one object contains
offload sections, only linker plugin and lto-wrapper can do it. But if linker
plugin is absent, the lto-wrapper have to open all objects, scan for all
sections, etc. Looks like too much overhead for a single diagnostic.
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-22 15:13 ` Thomas Schwinge
2016-02-22 18:08 ` Ilya Verbin
@ 2016-02-24 19:30 ` Ilya Verbin
2016-02-24 19:35 ` Jakub Jelinek
2016-05-10 15:41 ` [PATCH] Apply fix for PR68463 to RS6000 James Norris
2 siblings, 1 reply; 23+ messages in thread
From: Ilya Verbin @ 2016-02-24 19:30 UTC (permalink / raw)
To: Thomas Schwinge, Jakub Jelinek
Cc: bschmidt, gcc-patches, kirill.yukhin, Richard Biener, Tom de Vries
On Mon, Feb 22, 2016 at 16:13:07 +0100, Thomas Schwinge wrote:
> (..., and similar for others.) The if-exists spec function only works
> for absolute paths (I have not researched, why?), so it won't locate the
> files for relative -Bbuild-gcc/[...] prefixes, and linking will fail:
>
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x0): undefined reference to `__offload_func_table'
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x8): undefined reference to `__offload_funcs_end'
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x10): undefined reference to `__offload_var_table'
> /tmp/ccGajPD4.crtoffloadtable.o:(.rodata+0x18): undefined reference to `__offload_vars_end'
>
> If I use the absolute -B$PWD/build-gcc/[...], it works. (But there is no
> requirement for -B prefixes to be absolute, as far as I know.) Why not
> make it a hard error, though, if these files are missing? Can we use
> something like (untested pseudo-patch):
>
> +#ifdef ENABLE_OFFLOADING
> +# define CRTOFFLOADBEGIN "%{fopenacc|fopenmp:%:crtoffloadbegin%O%s}"
> +#else
> +# define CRTOFFLOADBEGIN ""
> +#endif
>
> @@ -49,14 +49,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 ")}"
Fixed. Actually ENABLE_OFFLOADING is always defined (to 0 or to 1).
> To the casual reader, skipping the first offload_files looks like a
> off-by-one error, so I suggest you add a comment "Skip the dummy item at
> the start of the list.", or similar.
Done.
> Ilya, then please remove
> libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims-2.c as part of
> your patch, unless Tom thinks it should be changed to a -flto test, but
> without -fno-use-linker-plugin?
Done.
Here is a follow up patch. OK for trunk? Bootstrapped and regtested.
Unfortunately I'm unable to run bootstrap-lto:
libdecnumber/dpd/decimal32.c:53:0: error: type of ‘decDigitsFromDPD’ does not match original declaration [-Werror=lto-type-mismatch]
[...]
diff --git a/gcc/config/gnu-user.h b/gcc/config/gnu-user.h
index 2fdb63c..b0bf40a 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
@@ -50,7 +58,7 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
%{fvtable-verify=none:%s; \
fvtable-verify=preinit:vtv_start_preinit.o%s; \
fvtable-verify=std:vtv_start.o%s} \
- %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
+ " CRTOFFLOADBEGIN
#else
#define GNU_USER_TARGET_STARTFILE_SPEC \
"%{!shared: %{pg|p|profile:gcrt1.o%s;:crt1.o%s}} \
@@ -58,7 +66,7 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
%{fvtable-verify=none:%s; \
fvtable-verify=preinit:vtv_start_preinit.o%s; \
fvtable-verify=std:vtv_start.o%s} \
- %{fopenacc|fopenmp:%:if-exists(crtoffloadbegin%O%s)}"
+ " CRTOFFLOADBEGIN
#endif
#undef STARTFILE_SPEC
#define STARTFILE_SPEC GNU_USER_TARGET_STARTFILE_SPEC
@@ -76,14 +84,14 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
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 \
- %{fopenacc|fopenmp:%:if-exists(crtoffloadend%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 \
- %{fopenacc|fopenmp:%:if-exists(crtoffloadend%O%s)}"
+ " CRTOFFLOADEND
#endif
#undef ENDFILE_SPEC
#define ENDFILE_SPEC GNU_USER_TARGET_ENDFILE_SPEC
diff --git a/libgcc/offloadstuff.c b/libgcc/offloadstuff.c
index a4ea3ac..4ab6397 100644
--- a/libgcc/offloadstuff.c
+++ b/libgcc/offloadstuff.c
@@ -40,7 +40,7 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
#include "tm.h"
#include "libgcc_tm.h"
-#if defined(HAVE_GAS_HIDDEN) && defined(ENABLE_OFFLOADING)
+#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"
diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c
index 35cb63a..51afc52 100644
--- a/lto-plugin/lto-plugin.c
+++ b/lto-plugin/lto-plugin.c
@@ -696,6 +696,7 @@ all_symbols_read_handler (void)
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)
{
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 eea8c7e..0000000
--- 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;
-}
-
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH][RFC][Offloading] Fix PR68463
2016-02-24 19:30 ` Ilya Verbin
@ 2016-02-24 19:35 ` Jakub Jelinek
0 siblings, 0 replies; 23+ messages in thread
From: Jakub Jelinek @ 2016-02-24 19:35 UTC (permalink / raw)
To: Ilya Verbin
Cc: Thomas Schwinge, bschmidt, gcc-patches, kirill.yukhin,
Richard Biener, Tom de Vries
On Wed, Feb 24, 2016 at 10:29:47PM +0300, Ilya Verbin wrote:
> Done.
> Here is a follow up patch. OK for trunk? Bootstrapped and regtested.
Ok with appropriate ChangeLog entry.
Jakub
^ permalink raw reply [flat|nested] 23+ messages in thread
* [PATCH] Apply fix for PR68463 to RS6000
@ 2016-05-10 15:41 ` James Norris
2016-05-10 15:49 ` David Edelsohn
2016-05-11 14:13 ` Thomas Schwinge
0 siblings, 2 replies; 23+ messages in thread
From: James Norris @ 2016-05-10 15:41 UTC (permalink / raw)
To: David Edelsohn; +Cc: Ilya Verbin, GCC Patches
[-- Attachment #1: Type: text/plain, Size: 581 bytes --]
Hi!
The fix for PR68463 (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=68463)
was missing code that prevented the fix from working on RS6000. The
attached patch adds the missing code for RS6000.
Bootstrapped and regtested on Power8.
OK for trunk?
Thanks!
Jim
=========
ChangeLog
* config/rs6000/sysv4.h (CRTOFFLOADBEGIN): Define. Add crtoffloadbegin.o
if offloading is enabled and -fopenacc or -fopenmp is specified.
(CRTOFFLOADEND): Likewise.
(STARTFILE_LINUX_SPEC): Add CRTOFFLOADBEGIN.
(ENDFILE_LINUX_SPEC): Add CRTOFFLOADEND.
[-- Attachment #2: power8.patch --]
[-- Type: text/x-patch, Size: 1866 bytes --]
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index dfbb1c3..1052980 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,11 @@
+2016-05-10 James Norris <jnorris@codesourcery.com>
+
+ * config/rs6000/sysv4.h (CRTOFFLOADBEGIN): Define. Add crtoffloadbegin.o
+ if offloading is enabled and -fopenacc or -fopenmp is specified.
+ (CRTOFFLOADEND): Likewise.
+ (STARTFILE_LINUX_SPEC): Add CRTOFFLOADBEGIN.
+ (ENDFILE_LINUX_SPEC): Add CRTOFFLOADEND.
+
2016-05-10 Jakub Jelinek <jakub@redhat.com>
PR target/70927
diff --git a/gcc/config/rs6000/sysv4.h b/gcc/config/rs6000/sysv4.h
index 46d2b4b..dbef378 100644
--- a/gcc/config/rs6000/sysv4.h
+++ b/gcc/config/rs6000/sysv4.h
@@ -749,21 +749,32 @@ ENDIAN_SELECT(" -mbig", " -mlittle", DEFAULT_ASM_ENDIAN)
%{!mnewlib: %{pthread:-lpthread} %{shared:-lc} \
%{!shared: %{profile:-lc_p} %{!profile:-lc}}}"
+#if ENABLE_OFFLOADING == 1
+#define CRTOFFLOADBEGIN "%{fopenacc|fopenmp:crtoffloadbegin%O%s}"
+#define CRTOFFLOADEND "%{fopenacc|fopenmp:crtoffloadend%O%s}"
+#else
+#define CRTOFFLOADBEGIN ""
+#define CRTOFFLOADEND ""
+#endif
+
#ifdef HAVE_LD_PIE
#define STARTFILE_LINUX_SPEC "\
%{!shared: %{pg|p|profile:gcrt1.o%s;pie:Scrt1.o%s;:crt1.o%s}} \
%{mnewlib:ecrti.o%s;:crti.o%s} \
-%{static:crtbeginT.o%s;shared|pie:crtbeginS.o%s;:crtbegin.o%s}"
+%{static:crtbeginT.o%s;shared|pie:crtbeginS.o%s;:crtbegin.o%s} \
+" CRTOFFLOADBEGIN
#else
#define STARTFILE_LINUX_SPEC "\
%{!shared: %{pg|p|profile:gcrt1.o%s;:crt1.o%s}} \
%{mnewlib:ecrti.o%s;:crti.o%s} \
-%{static:crtbeginT.o%s;shared|pie:crtbeginS.o%s;:crtbegin.o%s}"
+%{static:crtbeginT.o%s;shared|pie:crtbeginS.o%s;:crtbegin.o%s} \
+" CRTOFFLOADBEGIN
#endif
#define ENDFILE_LINUX_SPEC "\
%{shared|pie:crtendS.o%s;:crtend.o%s} \
-%{mnewlib:ecrtn.o%s;:crtn.o%s}"
+%{mnewlib:ecrtn.o%s;:crtn.o%s} \
+" CRTOFFLOADEND
#define LINK_START_LINUX_SPEC ""
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH] Apply fix for PR68463 to RS6000
2016-05-10 15:41 ` [PATCH] Apply fix for PR68463 to RS6000 James Norris
@ 2016-05-10 15:49 ` David Edelsohn
2016-05-10 15:59 ` Ilya Verbin
2016-05-11 14:13 ` Thomas Schwinge
1 sibling, 1 reply; 23+ messages in thread
From: David Edelsohn @ 2016-05-10 15:49 UTC (permalink / raw)
To: James Norris; +Cc: Ilya Verbin, GCC Patches, Segher Boessenkool
On Tue, May 10, 2016 at 11:39 AM, James Norris <jnorris@codesourcery.com> wrote:
> Hi!
>
> The fix for PR68463 (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=68463)
> was missing code that prevented the fix from working on RS6000. The
> attached patch adds the missing code for RS6000.
>
> Bootstrapped and regtested on Power8.
>
> OK for trunk?
>
> Thanks!
> Jim
>
> =========
>
> ChangeLog
>
> * config/rs6000/sysv4.h (CRTOFFLOADBEGIN): Define. Add
> crtoffloadbegin.o
> if offloading is enabled and -fopenacc or -fopenmp is specified.
> (CRTOFFLOADEND): Likewise.
> (STARTFILE_LINUX_SPEC): Add CRTOFFLOADBEGIN.
> (ENDFILE_LINUX_SPEC): Add CRTOFFLOADEND.
Why is this enabled for openmp? Not all openmp applications require offloading.
I see that the same logic is used in config/gnu-user.h, but I'm
curious about the need.
This is okay.
Thanks, David
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH] Apply fix for PR68463 to RS6000
2016-05-10 15:49 ` David Edelsohn
@ 2016-05-10 15:59 ` Ilya Verbin
0 siblings, 0 replies; 23+ messages in thread
From: Ilya Verbin @ 2016-05-10 15:59 UTC (permalink / raw)
To: David Edelsohn; +Cc: James Norris, GCC Patches, Segher Boessenkool
On Tue, May 10, 2016 at 11:48:53 -0400, David Edelsohn wrote:
> On Tue, May 10, 2016 at 11:39 AM, James Norris <jnorris@codesourcery.com> wrote:
> > The fix for PR68463 (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=68463)
> > was missing code that prevented the fix from working on RS6000. The
> > attached patch adds the missing code for RS6000.
> >
> > Bootstrapped and regtested on Power8.
> >
> > OK for trunk?
> >
> > Thanks!
> > Jim
> >
> > =========
> >
> > ChangeLog
> >
> > * config/rs6000/sysv4.h (CRTOFFLOADBEGIN): Define. Add
> > crtoffloadbegin.o
> > if offloading is enabled and -fopenacc or -fopenmp is specified.
> > (CRTOFFLOADEND): Likewise.
> > (STARTFILE_LINUX_SPEC): Add CRTOFFLOADBEGIN.
> > (ENDFILE_LINUX_SPEC): Add CRTOFFLOADEND.
>
> Why is this enabled for openmp? Not all openmp applications require offloading.
>
> I see that the same logic is used in config/gnu-user.h, but I'm
> curious about the need.
Yes, this adds a bit overhead to openmp applications without offloading (when
the compiler is configured with enabled offloading). But there is no way to
determine from the driver whether the application uses offloading or not, only
underlying lto-wrapper can determine this by analyzing object files.
-- Ilya
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH] Apply fix for PR68463 to RS6000
2016-05-10 15:41 ` [PATCH] Apply fix for PR68463 to RS6000 James Norris
2016-05-10 15:49 ` David Edelsohn
@ 2016-05-11 14:13 ` Thomas Schwinge
2016-05-11 14:18 ` James Norris
1 sibling, 1 reply; 23+ messages in thread
From: Thomas Schwinge @ 2016-05-11 14:13 UTC (permalink / raw)
To: James Norris, David Edelsohn; +Cc: Ilya Verbin, GCC Patches
[-- Attachment #1: Type: text/plain, Size: 1472 bytes --]
Hi!
On Tue, 10 May 2016 10:39:33 -0500, James Norris <jnorris@codesourcery.com> wrote:
> The fix for PR68463 (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=68463)
> was missing code that prevented the fix from working on RS6000. The
> attached patch adds the missing code for RS6000.
:-( Bah. When reviewing these changes,
<http://news.gmane.org/find-root.php?message_id=%3C8760xg7p8c.fsf%40hertz.schwinge.homeip.net%3E>,
I had been under the impression that all GNU/Linux targets use the
gcc/config/gnu-user.h file being patched there, hence my comment "I guess
we currently don't have to care about offloading configurations not using
the gnu-user.h file in which you modified the
STARTFILE_SPEC/ENDFILE_SPEC?" -- but as we now found out that hard way,
this PowerPC target does not actually, and so we need to repeat the
changes here:
> * config/rs6000/sysv4.h (CRTOFFLOADBEGIN): Define. Add crtoffloadbegin.o
> if offloading is enabled and -fopenacc or -fopenmp is specified.
> (CRTOFFLOADEND): Likewise.
> (STARTFILE_LINUX_SPEC): Add CRTOFFLOADBEGIN.
> (ENDFILE_LINUX_SPEC): Add CRTOFFLOADEND.
Should have added a "PR driver/68463" tag to the ChangeLog snippet, to
get this commit added to <https://gcc.gnu.org/PR68463>.
Are you also going to commit this to gcc-6-branch, where it is broken in
the very same way? You can do so "as obvious", without special approval.
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [PATCH] Apply fix for PR68463 to RS6000
2016-05-11 14:13 ` Thomas Schwinge
@ 2016-05-11 14:18 ` James Norris
0 siblings, 0 replies; 23+ messages in thread
From: James Norris @ 2016-05-11 14:18 UTC (permalink / raw)
To: Thomas Schwinge, David Edelsohn; +Cc: Ilya Verbin, GCC Patches
Thomas,
On 05/11/2016 09:13 AM, Thomas Schwinge wrote:
> Hi!
>
> On Tue, 10 May 2016 10:39:33 -0500, James Norris <jnorris@codesourcery.com> wrote:
>> The fix for PR68463 (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=68463)
>> was missing code that prevented the fix from working on RS6000. The
>> attached patch adds the missing code for RS6000.
>
> :-( Bah. When reviewing these changes,
> <http://news.gmane.org/find-root.php?message_id=%3C8760xg7p8c.fsf%40hertz.schwinge.homeip.net%3E>,
> I had been under the impression that all GNU/Linux targets use the
> gcc/config/gnu-user.h file being patched there, hence my comment "I guess
> we currently don't have to care about offloading configurations not using
> the gnu-user.h file in which you modified the
> STARTFILE_SPEC/ENDFILE_SPEC?" -- but as we now found out that hard way,
> this PowerPC target does not actually, and so we need to repeat the
> changes here:
>
>> * config/rs6000/sysv4.h (CRTOFFLOADBEGIN): Define. Add crtoffloadbegin.o
>> if offloading is enabled and -fopenacc or -fopenmp is specified.
>> (CRTOFFLOADEND): Likewise.
>> (STARTFILE_LINUX_SPEC): Add CRTOFFLOADBEGIN.
>> (ENDFILE_LINUX_SPEC): Add CRTOFFLOADEND.
>
> Should have added a "PR driver/68463" tag to the ChangeLog snippet, to
> get this commit added to <https://gcc.gnu.org/PR68463>.
Oppssss.
>
> Are you also going to commit this to gcc-6-branch, where it is broken in
> the very same way? You can do so "as obvious", without special approval.
Yes. And gomp4 too.
Jim
^ permalink raw reply [flat|nested] 23+ messages in thread
end of thread, other threads:[~2016-05-11 14:18 UTC | newest]
Thread overview: 23+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2016-01-14 21:26 [PATCH][RFC][Offloading] Fix PR68463 Ilya Verbin
2016-01-15 8:15 ` Richard Biener
2016-01-18 20:34 ` Ilya Verbin
2016-01-19 8:57 ` Richard Biener
2016-01-19 9:36 ` Jakub Jelinek
2016-01-19 13:32 ` Ilya Verbin
2016-02-10 17:20 ` Ilya Verbin
2016-02-19 14:53 ` Jakub Jelinek
2016-02-19 17:58 ` Mike Stump
2016-02-20 10:55 ` Ilya Verbin
2016-02-22 10:59 ` Jakub Jelinek
2016-02-22 15:13 ` Thomas Schwinge
2016-02-22 18:08 ` Ilya Verbin
2016-02-23 7:37 ` Tom de Vries
2016-02-24 16:13 ` Thomas Schwinge
2016-02-24 16:30 ` Ilya Verbin
2016-02-24 19:30 ` Ilya Verbin
2016-02-24 19:35 ` Jakub Jelinek
2016-05-10 15:41 ` [PATCH] Apply fix for PR68463 to RS6000 James Norris
2016-05-10 15:49 ` David Edelsohn
2016-05-10 15:59 ` Ilya Verbin
2016-05-11 14:13 ` Thomas Schwinge
2016-05-11 14:18 ` James Norris
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).