* [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper @ 2014-10-02 15:15 Ilya Verbin 2014-10-08 10:27 ` Jakub Jelinek ` (2 more replies) 0 siblings, 3 replies; 24+ messages in thread From: Ilya Verbin @ 2014-10-02 15:15 UTC (permalink / raw) To: Jakub Jelinek, Jan Hubicka, Richard Biener, gcc-patches Cc: Bernd Schmidt, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy Hello, With this patch lto-wrapper performs invocation of mkoffload tool for each offload target. This tool should be provided by the corresponding offload compiler. It will compile IR from .gnu.offload_lto_* sections into offload target code and embed the resultant code (offload image) into the new host's object file. Bootstrapped and regtested on top of patch 3. Is it OK for trunk? Thanks, -- Ilya 2014-10-02 Ilya Verbin <ilya.verbin@intel.com> Bernd Schmidt <bernds@codesourcery.com> Andrey Turetskiy <andrey.turetskiy@intel.com> Michael Zolotukhin <michael.v.zolotukhin@intel.com> gcc/ * gcc.c (spec_host_machine, accel_dir_suffix): New variables. (process_command): Tweak path construction for the possibility of being configured as an offload compiler. (main): Likewise. Look up specs in just_machine_suffix only if not ACCEL_COMPILER. Construct OFFLOAD_TARGET_NAMES environment variable if we have OFFLOAD_TARGETS. * lto-wrapper.c (OFFLOAD_TARGET_NAMES_ENV): Define. (offload_names, offloadbegin, offloadend): New static variables. (free_array_of_ptrs, parse_env_var, access_check, compile_offload_image) (compile_images_for_offload_targets, copy_file, find_offloadbeginend): New static functions. (run_gcc): Determine whether offload sections are present. If so, run compile_images_for_offload_targets and return the names of new generated objects to linker. If there are offload sections, but no LTO sections, then return the copies of input objects without link-time recompilation. lto-plugin/ * lto-plugin.c (OFFLOAD_SECTION, OFFLOAD_SECTION_LEN): Define. (struct plugin_objfile): Add new field "offload". (process_offload_section): New static function. (claim_file_handler): Claim file if it contains offload sections. --- diff --git a/gcc/gcc.c b/gcc/gcc.c index 47c4e28..82509a8 100644 --- a/gcc/gcc.c +++ b/gcc/gcc.c @@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION; /* The target machine. */ static const char *spec_machine = DEFAULT_TARGET_MACHINE; +static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE; /* Nonzero if cross-compiling. When -b is used, the value comes from the `specs' file. */ @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2 relative to the driver. */ static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX; +/* A prefix to be used when this is an accelerator compiler. */ +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX; + /* Subdirectory to use for locating libraries. Set by set_multilib_dir based on the compilation options. */ @@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count, } gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); - tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, + tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, dir_separator_str, NULL); /* Look for tools relative to the location from which the driver is running, or, if that is not available, the configured prefix. */ tooldir_prefix = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix, - spec_machine, dir_separator_str, - spec_version, dir_separator_str, tooldir_prefix2, NULL); + spec_host_machine, dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL); free (tooldir_prefix2); add_prefix (&exec_prefixes, @@ -6878,8 +6882,8 @@ main (int argc, char **argv) /* Read specs from a file if there is one. */ - machine_suffix = concat (spec_machine, dir_separator_str, - spec_version, dir_separator_str, NULL); + machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, NULL); just_machine_suffix = concat (spec_machine, dir_separator_str, NULL); specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true); @@ -6889,16 +6893,17 @@ main (int argc, char **argv) else init_spec (); +#ifndef ACCEL_COMPILER /* We need to check standard_exec_prefix/just_machine_suffix/specs for any override of as, ld and libraries. */ specs_file = (char *) alloca (strlen (standard_exec_prefix) + strlen (just_machine_suffix) + sizeof ("specs")); - strcpy (specs_file, standard_exec_prefix); strcat (specs_file, just_machine_suffix); strcat (specs_file, "specs"); if (access (specs_file, R_OK) == 0) read_specs (specs_file, true, false); +#endif /* Process any configure-time defaults specified for the command line options, via OPTION_DEFAULT_SPECS. */ @@ -7077,8 +7082,9 @@ main (int argc, char **argv) /* If we have a GCC_EXEC_PREFIX envvar, modify it for cpp's sake. */ if (gcc_exec_prefix) - gcc_exec_prefix = concat (gcc_exec_prefix, spec_machine, dir_separator_str, - spec_version, dir_separator_str, NULL); + gcc_exec_prefix = concat (gcc_exec_prefix, spec_host_machine, + dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, NULL); /* Now we have the specs. Set the `valid' bits for switches that match anything in any spec. */ @@ -7097,6 +7103,16 @@ main (int argc, char **argv) obstack_grow (&collect_obstack, argv[0], strlen (argv[0]) + 1); xputenv (XOBFINISH (&collect_obstack, char *)); + if (strlen (OFFLOAD_TARGETS) > 0) + { + obstack_init (&collect_obstack); + obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=", + sizeof ("OFFLOAD_TARGET_NAMES=") - 1); + obstack_grow (&collect_obstack, OFFLOAD_TARGETS, + strlen (OFFLOAD_TARGETS) + 1); + xputenv (XOBFINISH (&collect_obstack, char *)); + } + /* Set up to remember the pathname of the lto wrapper. */ if (have_c) diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c index 08fd090..2c9b503 100644 --- a/gcc/lto-wrapper.c +++ b/gcc/lto-wrapper.c @@ -49,6 +49,8 @@ along with GCC; see the file COPYING3. If not see #include "lto-section-names.h" #include "collect-utils.h" +#define OFFLOAD_TARGET_NAMES_ENV "OFFLOAD_TARGET_NAMES" + enum lto_mode_d { LTO_MODE_NONE, /* Not doing LTO. */ LTO_MODE_LTO, /* Normal LTO. */ @@ -63,6 +65,8 @@ static char *flto_out; static unsigned int nr; static char **input_names; static char **output_names; +static char **offload_names; +static const char *offloadbegin, *offloadend; static char *makefile; const char tool_name[] = "lto-wrapper"; @@ -358,6 +362,224 @@ merge_and_complain (struct cl_decoded_option **decoded_options, } } +/* Auxiliary function that frees elements of PTR and PTR itself. + N is number of elements to be freed. If PTR is NULL, nothing is freed. + If an element is NULL, subsequent elements are not freed. */ + +static void ** +free_array_of_ptrs (void **ptr, unsigned n) +{ + if (!ptr) + return NULL; + for (unsigned i = 0; i < n; i++) + { + if (!ptr[i]) + break; + free (ptr[i]); + } + free (ptr); + return NULL; +} + +/* Parse STR, saving found tokens into PVALUES and return their number. + Tokens are assumed to be delimited by ':'. If APPEND is non-null, + append it to every token we find. */ + +static unsigned +parse_env_var (const char *str, char ***pvalues, const char *append) +{ + const char *curval, *nextval; + char **values; + unsigned num = 1, i; + + curval = strchr (str, ':'); + while (curval) + { + num++; + curval = strchr (curval + 1, ':'); + } + + values = (char**) xmalloc (num * sizeof (char*)); + curval = str; + nextval = strchrnul (curval, ':'); + + int append_len = append ? strlen (append) : 0; + for (i = 0; i < num; i++) + { + int l = nextval - curval; + values[i] = (char*) xmalloc (l + 1 + append_len); + memcpy (values[i], curval, l); + values[i][l] = 0; + if (append) + strcat (values[i], append); + curval = nextval + 1; + nextval = strchrnul (curval, ':'); + } + *pvalues = values; + return num; +} + +/* Check whether NAME can be accessed in MODE. This is like access, + except that it never considers directories to be executable. */ + +static int +access_check (const char *name, int mode) +{ + if (mode == X_OK) + { + struct stat st; + + if (stat (name, &st) < 0 + || S_ISDIR (st.st_mode)) + return -1; + } + + return access (name, mode); +} + +/* Prepare a target image for offload TARGET, using mkoffload tool from + COMPILER_PATH. Return the name of the resultant object file. */ + +static char * +compile_offload_image (const char *target, const char *compiler_path, + unsigned in_argc, char *in_argv[]) +{ + char *filename = NULL; + char **argv; + char *suffix + = XALLOCAVEC (char, strlen ("/accel//mkoffload") + strlen (target) + 1); + strcpy (suffix, "/accel/"); + strcat (suffix, target); + strcat (suffix, "/mkoffload"); + + char **paths = NULL; + unsigned n_paths = parse_env_var (compiler_path, &paths, suffix); + + const char *compiler = NULL; + for (unsigned i = 0; i < n_paths; i++) + if (access_check (paths[i], X_OK) == 0) + { + compiler = paths[i]; + break; + } + + if (!compiler) + goto out; + + /* Generate temporary output file name. */ + filename = make_temp_file (".target.o"); + + struct obstack argv_obstack; + obstack_init (&argv_obstack); + obstack_ptr_grow (&argv_obstack, compiler); + obstack_ptr_grow (&argv_obstack, "-o"); + obstack_ptr_grow (&argv_obstack, filename); + + for (unsigned i = 1; i < in_argc; i++) + obstack_ptr_grow (&argv_obstack, in_argv[i]); + obstack_ptr_grow (&argv_obstack, NULL); + + argv = XOBFINISH (&argv_obstack, char **); + fork_execute (argv[0], argv, true); + obstack_free (&argv_obstack, NULL); + + out: + free_array_of_ptrs ((void **) paths, n_paths); + return filename; +} + + +/* The main routine dealing with offloading. + The routine builds a target image for each offload target. IN_ARGC and + IN_ARGV specify options and input object files. As all of them could contain + target sections, we pass them all to target compilers. */ + +static void +compile_images_for_offload_targets (unsigned in_argc, char *in_argv[]) +{ + char **names = NULL; + const char *target_names = getenv (OFFLOAD_TARGET_NAMES_ENV); + if (!target_names) + return; + unsigned num_targets = parse_env_var (target_names, &names, NULL); + + const char *compiler_path = getenv ("COMPILER_PATH"); + if (!compiler_path) + goto out; + + /* Prepare an image for each target and save the name of the resultant object + file to the OFFLOAD_NAMES array. It is terminated by a NULL entry. */ + offload_names = XCNEWVEC (char *, num_targets + 1); + for (unsigned i = 0; i < num_targets; i++) + { + offload_names[i] = compile_offload_image (names[i], compiler_path, + in_argc, in_argv); + if (!offload_names[i]) + fatal_error ("problem with building target image for %s\n", names[i]); + } + + out: + 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 ("reading input file"); + if (len > 0) + { + fwrite (buffer, 1, len, d); + if (ferror (d) != 0) + fatal_error ("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 ("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 ("installation error, can't find crtoffloadbegin.o"); + + free_array_of_ptrs ((void **) paths, n_paths); +} + /* Execute gcc. ARGC is the number of arguments. ARGV contains the arguments. */ static void @@ -378,6 +600,8 @@ run_gcc (unsigned argc, char *argv[]) unsigned int decoded_options_count; struct obstack argv_obstack; int new_head_argc; + bool have_lto = false; + bool have_offload = false; /* Get the driver and options. */ collect_gcc = getenv ("COLLECT_GCC"); @@ -426,6 +650,9 @@ run_gcc (unsigned argc, char *argv[]) close (fd); continue; } + if (simple_object_find_section (sobj, OFFLOAD_SECTION_NAME_PREFIX ".opts", + &offset, &length, &errmsg, &err)) + have_offload = true; if (!simple_object_find_section (sobj, LTO_SECTION_NAME_PREFIX "." "opts", &offset, &length, &errmsg, &err)) { @@ -433,6 +660,7 @@ run_gcc (unsigned argc, char *argv[]) close (fd); continue; } + have_lto = true; lseek (fd, file_offset + offset, SEEK_SET); data = (char *)xmalloc (length); read (fd, data, length); @@ -624,6 +852,54 @@ run_gcc (unsigned argc, char *argv[]) /* Remember at which point we can scrub args to re-use the commons. */ new_head_argc = obstack_object_size (&argv_obstack) / sizeof (void *); + if (have_offload) + { + compile_images_for_offload_targets (argc, argv); + 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)) + { + 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); + } + + /* By default linker does not discard .gnu.offload_lto_* sections. */ + const char *linker_script = make_temp_file ("_linker_script.x"); + FILE *stream = fopen (linker_script, "w"); + if (!stream) + fatal_error ("fopen %s: %m", linker_script); + fprintf (stream, "SECTIONS { /DISCARD/ : { *(" + OFFLOAD_SECTION_NAME_PREFIX "*) } }\n"); + fclose (stream); + printf ("%s\n", linker_script); + + goto finish; + } + if (lto_mode == LTO_MODE_LTO) { flto_out = make_temp_file (".lto.o"); @@ -850,6 +1126,10 @@ cont: obstack_free (&env_obstack, NULL); } + finish: + if (offloadend) + printf ("%s\n", offloadend); + obstack_free (&argv_obstack, NULL); } diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c index 910e23c..fb6555d 100644 --- a/lto-plugin/lto-plugin.c +++ b/lto-plugin/lto-plugin.c @@ -86,6 +86,8 @@ along with this program; see the file COPYING3. If not see #define LTO_SECTION_PREFIX ".gnu.lto_.symtab" #define LTO_SECTION_PREFIX_LEN (sizeof (LTO_SECTION_PREFIX) - 1) +#define OFFLOAD_SECTION ".gnu.offload_lto_.opts" +#define OFFLOAD_SECTION_LEN (sizeof (OFFLOAD_SECTION) - 1) /* The part of the symbol table the plugin has to keep track of. Note that we must keep SYMS until all_symbols_read is called to give the linker time to @@ -111,6 +113,7 @@ struct plugin_symtab struct plugin_objfile { int found; + int offload; simple_object_read *objfile; struct plugin_symtab *out; const struct ld_plugin_input_file *file; @@ -862,6 +865,21 @@ err: return 0; } +/* Find an offload section of an object file. */ + +static int +process_offload_section (void *data, const char *name, off_t offset, off_t len) +{ + if (!strncmp (name, OFFLOAD_SECTION, OFFLOAD_SECTION_LEN)) + { + struct plugin_objfile *obj = (struct plugin_objfile *) data; + obj->offload = 1; + return 0; + } + + return 1; +} + /* Callback used by gold to check if the plugin will claim FILE. Writes the result in CLAIMED. */ @@ -899,6 +917,7 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed) *claimed = 0; obj.file = file; obj.found = 0; + obj.offload = 0; obj.out = <o_file.symtab; errmsg = NULL; obj.objfile = simple_object_start_read (file->fd, file->offset, LTO_SEGMENT_NAME, @@ -920,7 +939,11 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed) goto err; } - if (obj.found == 0) + if (obj.objfile) + simple_object_find_sections (obj.objfile, process_offload_section, + &obj, &err); + + if (obj.found == 0 && obj.offload == 0) goto err; if (obj.found > 1) -- 1.7.1 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-02 15:15 [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Ilya Verbin @ 2014-10-08 10:27 ` Jakub Jelinek 2014-10-09 12:09 ` Ilya Verbin 2015-05-12 16:32 ` Thomas Schwinge 2016-02-19 19:42 ` [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Thomas Schwinge 2 siblings, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2014-10-08 10:27 UTC (permalink / raw) To: Ilya Verbin Cc: Jan Hubicka, Richard Biener, gcc-patches, Bernd Schmidt, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On Thu, Oct 02, 2014 at 07:14:57PM +0400, Ilya Verbin wrote: > @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2 > relative to the driver. */ > static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX; > > +/* A prefix to be used when this is an accelerator compiler. */ > +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX; Is ACCEL_DIR_SUFFIX here "" or something starting with "/ ? > @@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count, > } > > gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); > - tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, > + tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, > dir_separator_str, NULL); > > /* Look for tools relative to the location from which the driver is > running, or, if that is not available, the configured prefix. */ > tooldir_prefix > = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix, > - spec_machine, dir_separator_str, > - spec_version, dir_separator_str, tooldir_prefix2, NULL); > + spec_host_machine, dir_separator_str, spec_version, > + accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL); > free (tooldir_prefix2); > > add_prefix (&exec_prefixes, The reason I'm asking is that otherwise it seems gcc.c heavily uses dir_separator_str for the separators. Don't have any experience with targets where DIR_SEPARATOR is not / though, perhaps it is a non-issue. > @@ -6878,8 +6882,8 @@ main (int argc, char **argv) > > /* Read specs from a file if there is one. */ > > - machine_suffix = concat (spec_machine, dir_separator_str, > - spec_version, dir_separator_str, NULL); > + machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version, > + accel_dir_suffix, dir_separator_str, NULL); > just_machine_suffix = concat (spec_machine, dir_separator_str, NULL); > > specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true); > @@ -6889,16 +6893,17 @@ main (int argc, char **argv) > else > init_spec (); > > +#ifndef ACCEL_COMPILER > /* We need to check standard_exec_prefix/just_machine_suffix/specs > for any override of as, ld and libraries. */ > specs_file = (char *) alloca (strlen (standard_exec_prefix) > + strlen (just_machine_suffix) + sizeof ("specs")); > - > strcpy (specs_file, standard_exec_prefix); > strcat (specs_file, just_machine_suffix); > strcat (specs_file, "specs"); > if (access (specs_file, R_OK) == 0) > read_specs (specs_file, true, false); > +#endif Why do you want to disable specs reading for the accel compiler? Then users won't have the possibility to override defaults etc. easily... > @@ -7097,6 +7103,16 @@ main (int argc, char **argv) > obstack_grow (&collect_obstack, argv[0], strlen (argv[0]) + 1); > xputenv (XOBFINISH (&collect_obstack, char *)); > > + if (strlen (OFFLOAD_TARGETS) > 0) > + { > + obstack_init (&collect_obstack); > + obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=", > + sizeof ("OFFLOAD_TARGET_NAMES=") - 1); > + obstack_grow (&collect_obstack, OFFLOAD_TARGETS, > + strlen (OFFLOAD_TARGETS) + 1); > + xputenv (XOBFINISH (&collect_obstack, char *)); > + } > + I'm surprised to see the obstack_init call here, but looking at gcc.c, I'm surprised to see it in more places. I've always thought that obstack_init was something you invoke generally once on a given obstack object, then work with the obstack and then obstack_free (..., NULL) it at the end. Now, if it wants the obstack to be live until exit, it just would not obstack_free it. But calling obstack_init on the already initialized obstack results IMHO in memory leaks. It should be initialized just once somewhere. > /* Set up to remember the pathname of the lto wrapper. */ > > if (have_c) > diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c > index 08fd090..2c9b503 100644 > --- a/gcc/lto-wrapper.c > +++ b/gcc/lto-wrapper.c > @@ -49,6 +49,8 @@ along with GCC; see the file COPYING3. If not see > #include "lto-section-names.h" > #include "collect-utils.h" > > +#define OFFLOAD_TARGET_NAMES_ENV "OFFLOAD_TARGET_NAMES" Missing comment about the env var and what it is for. > +/* Prepare a target image for offload TARGET, using mkoffload tool from > + COMPILER_PATH. Return the name of the resultant object file. */ > + > +static char * > +compile_offload_image (const char *target, const char *compiler_path, > + unsigned in_argc, char *in_argv[]) > +{ > + char *filename = NULL; > + char **argv; > + char *suffix > + = XALLOCAVEC (char, strlen ("/accel//mkoffload") + strlen (target) + 1); Use sizeof ("/accel//mkoffload") + strlen (target) instead? > + strcpy (suffix, "/accel/"); > + strcat (suffix, target); > + strcat (suffix, "/mkoffload"); > + > + char **paths = NULL; > + unsigned n_paths = parse_env_var (compiler_path, &paths, suffix); > + > + const char *compiler = NULL; > + for (unsigned i = 0; i < n_paths; i++) > + if (access_check (paths[i], X_OK) == 0) > + { > + compiler = paths[i]; > + break; > + } > + > + if (!compiler) > + goto out; > + > + /* Generate temporary output file name. */ > + filename = make_temp_file (".target.o"); > + > + struct obstack argv_obstack; > + obstack_init (&argv_obstack); > + obstack_ptr_grow (&argv_obstack, compiler); > + obstack_ptr_grow (&argv_obstack, "-o"); > + obstack_ptr_grow (&argv_obstack, filename); > + > + for (unsigned i = 1; i < in_argc; i++) > + obstack_ptr_grow (&argv_obstack, in_argv[i]); > + obstack_ptr_grow (&argv_obstack, NULL); > + > + argv = XOBFINISH (&argv_obstack, char **); > + fork_execute (argv[0], argv, true); > + obstack_free (&argv_obstack, NULL); > + > + out: The goto is probably unnecessary here, indenting all the lines by 4 more spaces and using if (compiler) { ... } instead doesn't need any further warpping. > + free_array_of_ptrs ((void **) paths, n_paths); > + return filename; > +} > + > + /* By default linker does not discard .gnu.offload_lto_* sections. */ > + const char *linker_script = make_temp_file ("_linker_script.x"); > + FILE *stream = fopen (linker_script, "w"); > + if (!stream) > + fatal_error ("fopen %s: %m", linker_script); > + fprintf (stream, "SECTIONS { /DISCARD/ : { *(" > + OFFLOAD_SECTION_NAME_PREFIX "*) } }\n"); > + fclose (stream); > + printf ("%s\n", linker_script); > + > + goto finish; > + } Does this work with gold? Are there any other linkers that support plugins, but don't support linker scripts this way? Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-08 10:27 ` Jakub Jelinek @ 2014-10-09 12:09 ` Ilya Verbin 2014-10-09 12:13 ` Bernd Schmidt 2014-10-09 20:27 ` Ilya Verbin 0 siblings, 2 replies; 24+ messages in thread From: Ilya Verbin @ 2014-10-09 12:09 UTC (permalink / raw) To: Jakub Jelinek, Bernd Schmidt Cc: Jan Hubicka, Richard Biener, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On 08 Oct 12:26, Jakub Jelinek wrote: > On Thu, Oct 02, 2014 at 07:14:57PM +0400, Ilya Verbin wrote: > > @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2 > > relative to the driver. */ > > static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX; > > > > +/* A prefix to be used when this is an accelerator compiler. */ > > +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX; > > Is ACCEL_DIR_SUFFIX here "" or something starting with "/ ? > > The reason I'm asking is that otherwise it seems gcc.c heavily uses > dir_separator_str for the separators. Don't have any experience with > targets where DIR_SEPARATOR is not / though, perhaps it is a non-issue. It is "" for the host compiler and starts with "/" for the accel compiler. Looking at gcc/configure.ac, there are a lot of paths with slashes. So, I think it's ok to define ACCEL_DIR_SUFFIX in such a way. > > +#ifndef ACCEL_COMPILER > > /* We need to check standard_exec_prefix/just_machine_suffix/specs > > for any override of as, ld and libraries. */ > > specs_file = (char *) alloca (strlen (standard_exec_prefix) > > + strlen (just_machine_suffix) + sizeof ("specs")); > > - > > strcpy (specs_file, standard_exec_prefix); > > strcat (specs_file, just_machine_suffix); > > strcat (specs_file, "specs"); > > if (access (specs_file, R_OK) == 0) > > read_specs (specs_file, true, false); > > +#endif > > Why do you want to disable specs reading for the accel compiler? > Then users won't have the possibility to override defaults etc. easily... Bernd, Do you need this ifndef for PTX? Or I could remove it? > > @@ -7097,6 +7103,16 @@ main (int argc, char **argv) > > obstack_grow (&collect_obstack, argv[0], strlen (argv[0]) + 1); > > xputenv (XOBFINISH (&collect_obstack, char *)); > > > > + if (strlen (OFFLOAD_TARGETS) > 0) > > + { > > + obstack_init (&collect_obstack); > > + obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=", > > + sizeof ("OFFLOAD_TARGET_NAMES=") - 1); > > + obstack_grow (&collect_obstack, OFFLOAD_TARGETS, > > + strlen (OFFLOAD_TARGETS) + 1); > > + xputenv (XOBFINISH (&collect_obstack, char *)); > > + } > > + > > I'm surprised to see the obstack_init call here, but looking at > gcc.c, I'm surprised to see it in more places. > > I've always thought that obstack_init was something you invoke generally > once on a given obstack object, then work with the obstack and then > obstack_free (..., NULL) it at the end. Now, if it wants the obstack to be > live until exit, it just would not obstack_free it. But calling > obstack_init on the already initialized obstack results IMHO in memory > leaks. It should be initialized just once somewhere. Right, I removed obstack_init from this patch. > > +#define OFFLOAD_TARGET_NAMES_ENV "OFFLOAD_TARGET_NAMES" > > Missing comment about the env var and what it is for. Fixed. > > + char *suffix > > + = XALLOCAVEC (char, strlen ("/accel//mkoffload") + strlen (target) + 1); > > Use sizeof ("/accel//mkoffload") + strlen (target) instead? Done. > > + out: > > The goto is probably unnecessary here, indenting all the lines by > 4 more spaces and using if (compiler) { ... } instead doesn't need > any further warpping. Done. > > + /* By default linker does not discard .gnu.offload_lto_* sections. */ > > + const char *linker_script = make_temp_file ("_linker_script.x"); > > + FILE *stream = fopen (linker_script, "w"); > > + if (!stream) > > + fatal_error ("fopen %s: %m", linker_script); > > + fprintf (stream, "SECTIONS { /DISCARD/ : { *(" > > + OFFLOAD_SECTION_NAME_PREFIX "*) } }\n"); > > + fclose (stream); > > + printf ("%s\n", linker_script); > > + > > + goto finish; > > + } > > Does this work with gold? Are there any other linkers that support plugins, > but don't support linker scripts this way? Oops, gold does not support scripts, outputted from plugins :( "error: SECTIONS seen after other input files; try -T/--script" Probably, we should update default linker scripts in binutils? But without latest ld/gold all binaries compiled without -flto and with offloading will contain intermediate bytecode... Thanks, -- Ilya ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-09 12:09 ` Ilya Verbin @ 2014-10-09 12:13 ` Bernd Schmidt 2014-10-09 20:27 ` Ilya Verbin 1 sibling, 0 replies; 24+ messages in thread From: Bernd Schmidt @ 2014-10-09 12:13 UTC (permalink / raw) To: Ilya Verbin, Jakub Jelinek Cc: Jan Hubicka, Richard Biener, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On 10/09/2014 02:07 PM, Ilya Verbin wrote: >>> +#ifndef ACCEL_COMPILER >>> /* We need to check standard_exec_prefix/just_machine_suffix/specs >>> for any override of as, ld and libraries. */ >>> specs_file = (char *) alloca (strlen (standard_exec_prefix) >>> + strlen (just_machine_suffix) + sizeof ("specs")); >>> - >>> strcpy (specs_file, standard_exec_prefix); >>> strcat (specs_file, just_machine_suffix); >>> strcat (specs_file, "specs"); >>> if (access (specs_file, R_OK) == 0) >>> read_specs (specs_file, true, false); >>> +#endif >> >> Why do you want to disable specs reading for the accel compiler? >> Then users won't have the possibility to override defaults etc. easily... > > Bernd, > Do you need this ifndef for PTX? Or I could remove it? I suspect the paths aren't right. If that can be fixed it should be fine to remove the ifdef. Bernd ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-09 12:09 ` Ilya Verbin 2014-10-09 12:13 ` Bernd Schmidt @ 2014-10-09 20:27 ` Ilya Verbin 2014-10-10 7:13 ` Jakub Jelinek 1 sibling, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-09 20:27 UTC (permalink / raw) To: Jakub Jelinek, Richard Biener Cc: Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On 09 Oct 16:07, Ilya Verbin wrote: > > > + /* By default linker does not discard .gnu.offload_lto_* sections. */ > > > + const char *linker_script = make_temp_file ("_linker_script.x"); > > > + FILE *stream = fopen (linker_script, "w"); > > > + if (!stream) > > > + fatal_error ("fopen %s: %m", linker_script); > > > + fprintf (stream, "SECTIONS { /DISCARD/ : { *(" > > > + OFFLOAD_SECTION_NAME_PREFIX "*) } }\n"); > > > + fclose (stream); > > > + printf ("%s\n", linker_script); > > > + > > > + goto finish; > > > + } > > > > Does this work with gold? Are there any other linkers that support plugins, > > but don't support linker scripts this way? > > Oops, gold does not support scripts, outputted from plugins :( > "error: SECTIONS seen after other input files; try -T/--script" > > Probably, we should update default linker scripts in binutils? > But without latest ld/gold all binaries compiled without -flto and with > offloading will contain intermediate bytecode... Actually, this issue is not due to outputting a script from a plugin, gold just does not support partial linker scripts: https://sourceware.org/bugzilla/show_bug.cgi?id=17451 So it seems that discarding .gnu.offload_lto_* sections (like it is done for .gnu.lto_*) in the default ld and gold scripts is the right way? Thanks, -- Ilya ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-09 20:27 ` Ilya Verbin @ 2014-10-10 7:13 ` Jakub Jelinek 2014-10-10 16:52 ` Cary Coutant 0 siblings, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2014-10-10 7:13 UTC (permalink / raw) To: Ilya Verbin, Cary Coutant Cc: Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On Fri, Oct 10, 2014 at 12:07:03AM +0400, Ilya Verbin wrote: > On 09 Oct 16:07, Ilya Verbin wrote: > > > > + /* By default linker does not discard .gnu.offload_lto_* sections. */ > > > > + const char *linker_script = make_temp_file ("_linker_script.x"); > > > > + FILE *stream = fopen (linker_script, "w"); > > > > + if (!stream) > > > > + fatal_error ("fopen %s: %m", linker_script); > > > > + fprintf (stream, "SECTIONS { /DISCARD/ : { *(" > > > > + OFFLOAD_SECTION_NAME_PREFIX "*) } }\n"); > > > > + fclose (stream); > > > > + printf ("%s\n", linker_script); > > > > + > > > > + goto finish; > > > > + } > > > > > > Does this work with gold? Are there any other linkers that support plugins, > > > but don't support linker scripts this way? > > > > Oops, gold does not support scripts, outputted from plugins :( > > "error: SECTIONS seen after other input files; try -T/--script" > > > > Probably, we should update default linker scripts in binutils? > > But without latest ld/gold all binaries compiled without -flto and with > > offloading will contain intermediate bytecode... > > Actually, this issue is not due to outputting a script from a plugin, > gold just does not support partial linker scripts: > https://sourceware.org/bugzilla/show_bug.cgi?id=17451 > > So it seems that discarding .gnu.offload_lto_* sections (like it is done for > .gnu.lto_*) in the default ld and gold scripts is the right way? I must say I'm not very much familiar with the linker plugin API, but it surprises me that discarding sections is not something it allows. Anyway, can you do the partial linker script for the bfd linker (is there a way to determine from the linker plugin API if it is gold or bfd ld?), and for gold for the time being perhaps strip the sections in lto-wrapper? and feed the ET_REL objects with the sections stripped back to the linker through the plugin API? Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-10 7:13 ` Jakub Jelinek @ 2014-10-10 16:52 ` Cary Coutant 2014-10-10 17:01 ` Jakub Jelinek 0 siblings, 1 reply; 24+ messages in thread From: Cary Coutant @ 2014-10-10 16:52 UTC (permalink / raw) To: Jakub Jelinek Cc: Ilya Verbin, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy The linker already has a --strip-lto-sections option, and it's on by default. I'll approve a patch that modifies gold to recognize .gnu.offload_lto.* sections as part of --strip-lto-sections. Really, though, you should be setting the SHF_EXCLUDE bit on these sections. Do that and no special-casing will be necessary. Generating a linker script on the fly to discard these sections is, to me, rather hacky. There are better ways to do it. -cary On Thu, Oct 9, 2014 at 11:53 PM, Jakub Jelinek <jakub@redhat.com> wrote: > On Fri, Oct 10, 2014 at 12:07:03AM +0400, Ilya Verbin wrote: >> On 09 Oct 16:07, Ilya Verbin wrote: >> > > > + /* By default linker does not discard .gnu.offload_lto_* sections. */ >> > > > + const char *linker_script = make_temp_file ("_linker_script.x"); >> > > > + FILE *stream = fopen (linker_script, "w"); >> > > > + if (!stream) >> > > > + fatal_error ("fopen %s: %m", linker_script); >> > > > + fprintf (stream, "SECTIONS { /DISCARD/ : { *(" >> > > > + OFFLOAD_SECTION_NAME_PREFIX "*) } }\n"); >> > > > + fclose (stream); >> > > > + printf ("%s\n", linker_script); >> > > > + >> > > > + goto finish; >> > > > + } >> > > >> > > Does this work with gold? Are there any other linkers that support plugins, >> > > but don't support linker scripts this way? >> > >> > Oops, gold does not support scripts, outputted from plugins :( >> > "error: SECTIONS seen after other input files; try -T/--script" >> > >> > Probably, we should update default linker scripts in binutils? >> > But without latest ld/gold all binaries compiled without -flto and with >> > offloading will contain intermediate bytecode... >> >> Actually, this issue is not due to outputting a script from a plugin, >> gold just does not support partial linker scripts: >> https://sourceware.org/bugzilla/show_bug.cgi?id=17451 >> >> So it seems that discarding .gnu.offload_lto_* sections (like it is done for >> .gnu.lto_*) in the default ld and gold scripts is the right way? > > I must say I'm not very much familiar with the linker plugin API, but it > surprises me that discarding sections is not something it allows. > Anyway, can you do the partial linker script for the bfd linker (is there > a way to determine from the linker plugin API if it is gold or bfd ld?), and > for gold for the time being perhaps strip the sections in lto-wrapper? and > feed the ET_REL objects with the sections stripped back to the linker > through the plugin API? > > Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-10 16:52 ` Cary Coutant @ 2014-10-10 17:01 ` Jakub Jelinek 2014-10-10 17:10 ` Cary Coutant 2014-10-13 22:47 ` Ilya Verbin 0 siblings, 2 replies; 24+ messages in thread From: Jakub Jelinek @ 2014-10-10 17:01 UTC (permalink / raw) To: Cary Coutant Cc: Ilya Verbin, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On Fri, Oct 10, 2014 at 09:51:02AM -0700, Cary Coutant wrote: > The linker already has a --strip-lto-sections option, and it's on by > default. I'll approve a patch that modifies gold to recognize > .gnu.offload_lto.* sections as part of --strip-lto-sections. > > Really, though, you should be setting the SHF_EXCLUDE bit on these > sections. Do that and no special-casing will be necessary. For that I guess lhd_begin_section would need to replace: section = get_section (name, SECTION_DEBUG, NULL); with: section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL); either just for the .gnu.offload_lto prefixed section, or all. The question is what will old assemblers and/or linkers do with that, and if there are any that support linker plugins, but not SHF_EXCLUDE. Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-10 17:01 ` Jakub Jelinek @ 2014-10-10 17:10 ` Cary Coutant 2014-10-13 22:47 ` Ilya Verbin 1 sibling, 0 replies; 24+ messages in thread From: Cary Coutant @ 2014-10-10 17:10 UTC (permalink / raw) To: Jakub Jelinek Cc: Ilya Verbin, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy > The question is what will old assemblers and/or linkers do with that, and > if there are any that support linker plugins, but not SHF_EXCLUDE. If it helps answer that question, SHF_EXCLUDE support has been in gold for 6 years, and in gas for 4. -cary ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-10 17:01 ` Jakub Jelinek 2014-10-10 17:10 ` Cary Coutant @ 2014-10-13 22:47 ` Ilya Verbin 2014-10-14 9:41 ` Jakub Jelinek 1 sibling, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-13 22:47 UTC (permalink / raw) To: Jakub Jelinek Cc: Cary Coutant, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On 10 Oct 19:00, Jakub Jelinek wrote: > On Fri, Oct 10, 2014 at 09:51:02AM -0700, Cary Coutant wrote: > > The linker already has a --strip-lto-sections option, and it's on by > > default. I'll approve a patch that modifies gold to recognize > > .gnu.offload_lto.* sections as part of --strip-lto-sections. > > > > Really, though, you should be setting the SHF_EXCLUDE bit on these > > sections. Do that and no special-casing will be necessary. > > For that I guess > lhd_begin_section > would need to replace: > section = get_section (name, SECTION_DEBUG, NULL); > with: > section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL); > either just for the .gnu.offload_lto prefixed section, or all. > The question is what will old assemblers and/or linkers do with that, and > if there are any that support linker plugins, but not SHF_EXCLUDE. I've tried to set SECTION_EXCLUDE bit with as+ld version 2.20.51 and got a lot of warnings like: /tmp/ccg7P7iS.s:2: Warning: entity size for SHF_MERGE not specified /tmp/ccg7P7iS.s:2: Warning: group name for SHF_GROUP not specified as: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166' /usr/bin/ld: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166' I think, it can be placed under such ifdef: #if defined (HAVE_SECTION_EXCLUDE) && HAVE_SECTION_EXCLUDE == 1 section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL); #else section = get_section (name, SECTION_DEBUG, NULL); #endif Currently there is HAVE_GAS_SECTION_EXCLUDE implemented in gcc/configure.ac, and HAVE_SECTION_EXCLUDE can use it + check a version of the linker. As for old assemblers and/or linkers, which doesn't support SHF_EXCLUDE: On 10 Oct 08:53, Jakub Jelinek wrote: > Anyway, can you do the partial linker script for the bfd linker (is there > a way to determine from the linker plugin API if it is gold or bfd ld?), and > for gold for the time being perhaps strip the sections in lto-wrapper? and > feed the ET_REL objects with the sections stripped back to the linker > through the plugin API? Yes, it's possible to determine the versions of gold and ld.bfd from the plugin (plugin-api.h:ld_plugin_tag). But the problem is not only with plugins. If linker plugin in disabled, collect2 runs lto-wrapper directly to compile offload_lto sections and/or to perform usual LTO: https://gcc.gnu.org/wiki/Offloading#Compilation_without_-flto_and_without_linker_plugin And if -flto is absent, but offload_lto sections were emitted, the final binary will contain them. Perhaps, collect2 could execute and parse `ld --version` to determine whether it supports SHF_EXCLUDE. Then, we need somehow pass this information from plugin or collect2 to lto-wrapper. Maybe using an option like --strip-offload-sections, but currently lto-wrapper doesn't have its own options. And finally, under this option lto-wrapper will execute `objcopy --remove-section=<offload_lto>` and return stripped objects back to the linker or to collect2. So, should I implement this stuff, or maybe just leave offload_lto sections in the case of old binutils? :) Thanks, -- Ilya ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-13 22:47 ` Ilya Verbin @ 2014-10-14 9:41 ` Jakub Jelinek 2014-10-15 14:27 ` Ilya Verbin 2014-10-15 17:03 ` Cary Coutant 0 siblings, 2 replies; 24+ messages in thread From: Jakub Jelinek @ 2014-10-14 9:41 UTC (permalink / raw) To: Ilya Verbin Cc: Cary Coutant, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On Tue, Oct 14, 2014 at 02:42:47AM +0400, Ilya Verbin wrote: > > For that I guess > > lhd_begin_section > > would need to replace: > > section = get_section (name, SECTION_DEBUG, NULL); > > with: > > section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL); > > either just for the .gnu.offload_lto prefixed section, or all. > > The question is what will old assemblers and/or linkers do with that, and > > if there are any that support linker plugins, but not SHF_EXCLUDE. > > I've tried to set SECTION_EXCLUDE bit with as+ld version 2.20.51 and got a lot > of warnings like: > > /tmp/ccg7P7iS.s:2: Warning: entity size for SHF_MERGE not specified > /tmp/ccg7P7iS.s:2: Warning: group name for SHF_GROUP not specified > as: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166' > /usr/bin/ld: /tmp/ccKFKXfc.o: warning: sh_link not set for section `.gnu.lto_main.11d9780ff2ebf166' > > I think, it can be placed under such ifdef: > > #if defined (HAVE_SECTION_EXCLUDE) && HAVE_SECTION_EXCLUDE == 1 > section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL); > #else > section = get_section (name, SECTION_DEBUG, NULL); > #endif > > Currently there is HAVE_GAS_SECTION_EXCLUDE implemented in gcc/configure.ac, and > HAVE_SECTION_EXCLUDE can use it + check a version of the linker. My preference would be to add the | SECTION_EXCLUDE unconditionally, and instead guard the if (flags & SECTION_EXCLUDE) *f++ = 'e'; in varasm.c (default_elf_asm_named_section). The only other user of SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change ok with you? If you have new gas and old linker, I'd expect it would just ignore SHF_EXCLUDE. Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-14 9:41 ` Jakub Jelinek @ 2014-10-15 14:27 ` Ilya Verbin 2014-10-15 14:46 ` Jakub Jelinek 2014-10-15 17:03 ` Cary Coutant 1 sibling, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-15 14:27 UTC (permalink / raw) To: Jakub Jelinek, Bernd Schmidt Cc: Cary Coutant, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On 14 Oct 11:40, Jakub Jelinek wrote: > My preference would be to add the | SECTION_EXCLUDE unconditionally, and > instead guard the > if (flags & SECTION_EXCLUDE) > *f++ = 'e'; > in varasm.c (default_elf_asm_named_section). The only other user of > SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change > ok with you? Done. But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac: gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],, [.section foo1,"e" .byte 0,0,0,0]) does not work properly. Maybe it works on "cygwin* | pe | mingw32* | interix*" targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude sections) it successfully assembles conftest.s into conftest.o (with warnings) and HAVE_GAS_SECTION_EXCLUDE becomes defined. $ cat conftest.s .section foo1,"e" .byte 0,0,0,0 $ as conftest.s -o conftest.o conftest.s: Assembler messages: conftest.s:1: Warning: entity size for SHF_MERGE not specified conftest.s:1: Warning: group name for SHF_GROUP not specified as: a.out: warning: sh_link not set for section `foo1' $ readelf -t conftest.o [4] foo1 PROGBITS 0000000000000000 0000000000000040 0 0000000000000004 0000000000000000 0 1 [fffffffffffffdef]: WRITE, ALLOC, EXEC, STRINGS, INFO LINK, LINK ORDER, OS NONCONF, TLS, ORDERED, EXCLUDE, OS (000000000ff00000), PROC (0000000030000000), UNKNOWN (ffffffff000ff808) And therefore instead of assembling "e" flag, I had to check the version of gas. If it is >= 2.21.51 then HAVE_GAS_SECTION_EXCLUDE becomes defined. > If you have new gas and old linker, I'd expect it would just ignore > SHF_EXCLUDE. Yes, gas 2.21.51 + ld 2.20.51 build sections with SHF_EXCLUDE bit without warnings, ld just ignores it. I've updated the patch. It also fixes the issue with paths to accel specs. On 09 Oct 14:08, Bernd Schmidt wrote: > >>Why do you want to disable specs reading for the accel compiler? > >>Then users won't have the possibility to override defaults etc. easily... > I suspect the paths aren't right. If that can be fixed it should be > fine to remove the ifdef. Bootstrapped and regtested on i686-linux and x86_64-linux. Is it OK now? Thanks, -- Ilya gcc/ * configure: Regenerate. * configure.ac (HAVE_GAS_SECTION_EXCLUDE): Define if GNU assembler version >= 2.21.51 is used. * gcc.c (spec_host_machine, accel_dir_suffix): New variables. (process_command): Tweak path construction for the possibility of being configured as an offload compiler. (driver::maybe_putenv_OFFLOAD_TARGETS): New function. (driver::main): Call maybe_putenv_OFFLOAD_TARGETS. (driver::set_up_specs): Tweak path construction for the possibility of being configured as an offload compiler. * langhooks.c (lhd_begin_section): Set SECTION_EXCLUDE flag. * lto-wrapper.c (OFFLOAD_TARGET_NAMES_ENV): Define. (offload_names, offloadbegin, offloadend): New static variables. (free_array_of_ptrs, parse_env_var, access_check, compile_offload_image) (compile_images_for_offload_targets, copy_file, find_offloadbeginend): New static functions. (run_gcc): Determine whether offload sections are present. If so, run compile_images_for_offload_targets and return the names of new generated objects to linker. If there are offload sections, but no LTO sections, then return the copies of input objects without link-time recompilation. * varasm.c (default_elf_asm_named_section): Guard SECTION_EXCLUDE with ifdef HAVE_GAS_SECTION_EXCLUDE. lto-plugin/ * lto-plugin.c (OFFLOAD_SECTION, OFFLOAD_SECTION_LEN): Define. (struct plugin_objfile): Add new field "offload". (process_offload_section): New static function. (claim_file_handler): Claim file if it contains offload sections. --- diff --git a/gcc/configure b/gcc/configure index ff1e398..5e61ac8 100755 --- a/gcc/configure +++ b/gcc/configure @@ -24755,13 +24755,37 @@ fi { $as_echo "$as_me:${as_lineno-$LINENO}: result: $gcc_cv_as_section_has_e" >&5 $as_echo "$gcc_cv_as_section_has_e" >&6; } + ;; + esac + + # GAS versions before 2.21.51 do not support the section flag 'e' for + # specifying an excluded section. + if test "$gcc_cv_as_section_has_e" != "yes"; then + as_ver=`$gcc_cv_as --version 2>/dev/null | sed 1q` + if echo "$as_ver" | grep GNU > /dev/null; then + as_vers=`echo $as_ver | sed -n \ + -e 's,^.*[ ]\([0-9][0-9]*\.[0-9][0-9]*.*\)$,\1,p'` + as_major=`expr "$as_vers" : '\([0-9]*\)'` + as_minor=`expr "$as_vers" : '[0-9]*\.\([0-9]*\)'` + as_patch=`expr "$as_vers" : '[0-9]*\.[0-9]*\.\([0-9]*\)'` + if test $as_major -ge 3; then + gcc_cv_as_section_has_e=yes + elif test $as_major -eq 2; then + if test $as_minor -ge 22; then + gcc_cv_as_section_has_e=yes + elif test $as_minor -eq 21; then + if test $as_patch -ge 51; then + gcc_cv_as_section_has_e=yes + fi + fi + fi + fi + fi cat >>confdefs.h <<_ACEOF #define HAVE_GAS_SECTION_EXCLUDE `if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi` _ACEOF - ;; - esac { $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for filds and fists mnemonics" >&5 $as_echo_n "checking assembler for filds and fists mnemonics... " >&6; } diff --git a/gcc/configure.ac b/gcc/configure.ac index 05a55f4..5edbebe 100644 --- a/gcc/configure.ac +++ b/gcc/configure.ac @@ -3843,12 +3843,36 @@ foo: nop [2,22,51],, [.section foo1,"e" .byte 0,0,0,0]) - AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE, - [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`], - [Define if your assembler supports specifying the section flag e.]) ;; esac + # GAS versions before 2.21.51 do not support the section flag 'e' for + # specifying an excluded section. + [if test "$gcc_cv_as_section_has_e" != "yes"; then + as_ver=`$gcc_cv_as --version 2>/dev/null | sed 1q` + if echo "$as_ver" | grep GNU > /dev/null; then + as_vers=`echo $as_ver | sed -n \ + -e 's,^.*[ ]\([0-9][0-9]*\.[0-9][0-9]*.*\)$,\1,p'` + as_major=`expr "$as_vers" : '\([0-9]*\)'` + as_minor=`expr "$as_vers" : '[0-9]*\.\([0-9]*\)'` + as_patch=`expr "$as_vers" : '[0-9]*\.[0-9]*\.\([0-9]*\)'` + if test $as_major -ge 3; then + gcc_cv_as_section_has_e=yes + elif test $as_major -eq 2; then + if test $as_minor -ge 22; then + gcc_cv_as_section_has_e=yes + elif test $as_minor -eq 21; then + if test $as_patch -ge 51; then + gcc_cv_as_section_has_e=yes + fi + fi + fi + fi + fi] + AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE, + [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`], + [Define if your assembler supports specifying the section flag e.]) + gcc_GAS_CHECK_FEATURE([filds and fists mnemonics], gcc_cv_as_ix86_filds,,, [filds mem; fists mem],, diff --git a/gcc/gcc.c b/gcc/gcc.c index 71c76f8..4619fe7 100644 --- a/gcc/gcc.c +++ b/gcc/gcc.c @@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION; /* The target machine. */ static const char *spec_machine = DEFAULT_TARGET_MACHINE; +static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE; /* Nonzero if cross-compiling. When -b is used, the value comes from the `specs' file. */ @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2 relative to the driver. */ static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX; +/* A prefix to be used when this is an accelerator compiler. */ +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX; + /* Subdirectory to use for locating libraries. Set by set_multilib_dir based on the compilation options. */ @@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count, } gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); - tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, + tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, dir_separator_str, NULL); /* Look for tools relative to the location from which the driver is running, or, if that is not available, the configured prefix. */ tooldir_prefix = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix, - spec_machine, dir_separator_str, - spec_version, dir_separator_str, tooldir_prefix2, NULL); + spec_host_machine, dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL); free (tooldir_prefix2); add_prefix (&exec_prefixes, @@ -6742,6 +6746,7 @@ class driver void set_up_specs () const; void putenv_COLLECT_GCC (const char *argv0) const; void maybe_putenv_COLLECT_LTO_WRAPPER () const; + void maybe_putenv_OFFLOAD_TARGETS () const; void handle_unrecognized_options () const; int maybe_print_and_exit () const; bool prepare_infiles (); @@ -6784,6 +6789,7 @@ driver::main (int argc, char **argv) set_up_specs (); putenv_COLLECT_GCC (argv[0]); maybe_putenv_COLLECT_LTO_WRAPPER (); + maybe_putenv_OFFLOAD_TARGETS (); handle_unrecognized_options (); if (!maybe_print_and_exit ()) @@ -6953,6 +6959,7 @@ driver::build_multilib_strings () const void driver::set_up_specs () const { + const char *spec_machine_suffix; char *specs_file; size_t i; @@ -6976,8 +6983,8 @@ driver::set_up_specs () const /* Read specs from a file if there is one. */ - machine_suffix = concat (spec_machine, dir_separator_str, - spec_version, dir_separator_str, NULL); + machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, NULL); just_machine_suffix = concat (spec_machine, dir_separator_str, NULL); specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true); @@ -6987,13 +6994,18 @@ driver::set_up_specs () const else init_spec (); - /* We need to check standard_exec_prefix/just_machine_suffix/specs +#ifdef ACCEL_COMPILER + spec_machine_suffix = machine_suffix; +#else + spec_machine_suffix = just_machine_suffix; +#endif + + /* We need to check standard_exec_prefix/spec_machine_suffix/specs for any override of as, ld and libraries. */ specs_file = (char *) alloca (strlen (standard_exec_prefix) - + strlen (just_machine_suffix) + sizeof ("specs")); - + + strlen (spec_machine_suffix) + sizeof ("specs")); strcpy (specs_file, standard_exec_prefix); - strcat (specs_file, just_machine_suffix); + strcat (specs_file, spec_machine_suffix); strcat (specs_file, "specs"); if (access (specs_file, R_OK) == 0) read_specs (specs_file, true, false); @@ -7175,8 +7187,9 @@ driver::set_up_specs () const /* If we have a GCC_EXEC_PREFIX envvar, modify it for cpp's sake. */ if (gcc_exec_prefix) - gcc_exec_prefix = concat (gcc_exec_prefix, spec_machine, dir_separator_str, - spec_version, dir_separator_str, NULL); + gcc_exec_prefix = concat (gcc_exec_prefix, spec_host_machine, + dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, NULL); /* Now we have the specs. Set the `valid' bits for switches that match anything in any spec. */ @@ -7227,6 +7240,21 @@ driver::maybe_putenv_COLLECT_LTO_WRAPPER () const } +/* Set up to remember the names of offload targets. */ + +void +driver::maybe_putenv_OFFLOAD_TARGETS () const +{ + if (strlen (OFFLOAD_TARGETS) > 0) + { + obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=", + sizeof ("OFFLOAD_TARGET_NAMES=") - 1); + obstack_grow (&collect_obstack, OFFLOAD_TARGETS, + strlen (OFFLOAD_TARGETS) + 1); + xputenv (XOBFINISH (&collect_obstack, char *)); + } +} + /* Reject switches that no pass was interested in. */ void diff --git a/gcc/langhooks.c b/gcc/langhooks.c index 7d4c294..4bdeaa0 100644 --- a/gcc/langhooks.c +++ b/gcc/langhooks.c @@ -660,7 +660,7 @@ lhd_begin_section (const char *name) saved_section = text_section; /* Create a new section and switch to it. */ - section = get_section (name, SECTION_DEBUG, NULL); + section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL); switch_to_section (section); } diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c index 8033b15..cbda36b 100644 --- a/gcc/lto-wrapper.c +++ b/gcc/lto-wrapper.c @@ -49,6 +49,10 @@ along with GCC; see the file COPYING3. If not see #include "lto-section-names.h" #include "collect-utils.h" +/* Environment variable, used for passing the names of offload targets from GCC + driver to lto-wrapper. */ +#define OFFLOAD_TARGET_NAMES_ENV "OFFLOAD_TARGET_NAMES" + enum lto_mode_d { LTO_MODE_NONE, /* Not doing LTO. */ LTO_MODE_LTO, /* Normal LTO. */ @@ -63,6 +67,8 @@ static char *flto_out; static unsigned int nr; static char **input_names; static char **output_names; +static char **offload_names; +static const char *offloadbegin, *offloadend; static char *makefile; const char tool_name[] = "lto-wrapper"; @@ -364,6 +370,223 @@ merge_and_complain (struct cl_decoded_option **decoded_options, } } +/* Auxiliary function that frees elements of PTR and PTR itself. + N is number of elements to be freed. If PTR is NULL, nothing is freed. + If an element is NULL, subsequent elements are not freed. */ + +static void ** +free_array_of_ptrs (void **ptr, unsigned n) +{ + if (!ptr) + return NULL; + for (unsigned i = 0; i < n; i++) + { + if (!ptr[i]) + break; + free (ptr[i]); + } + free (ptr); + return NULL; +} + +/* Parse STR, saving found tokens into PVALUES and return their number. + Tokens are assumed to be delimited by ':'. If APPEND is non-null, + append it to every token we find. */ + +static unsigned +parse_env_var (const char *str, char ***pvalues, const char *append) +{ + const char *curval, *nextval; + char **values; + unsigned num = 1, i; + + curval = strchr (str, ':'); + while (curval) + { + num++; + curval = strchr (curval + 1, ':'); + } + + values = (char**) xmalloc (num * sizeof (char*)); + curval = str; + nextval = strchrnul (curval, ':'); + + int append_len = append ? strlen (append) : 0; + for (i = 0; i < num; i++) + { + int l = nextval - curval; + values[i] = (char*) xmalloc (l + 1 + append_len); + memcpy (values[i], curval, l); + values[i][l] = 0; + if (append) + strcat (values[i], append); + curval = nextval + 1; + nextval = strchrnul (curval, ':'); + } + *pvalues = values; + return num; +} + +/* Check whether NAME can be accessed in MODE. This is like access, + except that it never considers directories to be executable. */ + +static int +access_check (const char *name, int mode) +{ + if (mode == X_OK) + { + struct stat st; + + if (stat (name, &st) < 0 + || S_ISDIR (st.st_mode)) + return -1; + } + + return access (name, mode); +} + +/* Prepare a target image for offload TARGET, using mkoffload tool from + COMPILER_PATH. Return the name of the resultant object file. */ + +static char * +compile_offload_image (const char *target, const char *compiler_path, + unsigned in_argc, char *in_argv[]) +{ + char *filename = NULL; + char **argv; + char *suffix + = XALLOCAVEC (char, sizeof ("/accel//mkoffload") + strlen (target)); + strcpy (suffix, "/accel/"); + strcat (suffix, target); + strcat (suffix, "/mkoffload"); + + char **paths = NULL; + unsigned n_paths = parse_env_var (compiler_path, &paths, suffix); + + const char *compiler = NULL; + for (unsigned i = 0; i < n_paths; i++) + if (access_check (paths[i], X_OK) == 0) + { + compiler = paths[i]; + break; + } + + if (compiler) + { + /* Generate temporary output file name. */ + filename = make_temp_file (".target.o"); + + struct obstack argv_obstack; + obstack_init (&argv_obstack); + obstack_ptr_grow (&argv_obstack, compiler); + obstack_ptr_grow (&argv_obstack, "-o"); + obstack_ptr_grow (&argv_obstack, filename); + + for (unsigned i = 1; i < in_argc; i++) + obstack_ptr_grow (&argv_obstack, in_argv[i]); + obstack_ptr_grow (&argv_obstack, NULL); + + argv = XOBFINISH (&argv_obstack, char **); + fork_execute (argv[0], argv, true); + obstack_free (&argv_obstack, NULL); + } + + free_array_of_ptrs ((void **) paths, n_paths); + return filename; +} + + +/* The main routine dealing with offloading. + The routine builds a target image for each offload target. IN_ARGC and + IN_ARGV specify options and input object files. As all of them could contain + target sections, we pass them all to target compilers. */ + +static void +compile_images_for_offload_targets (unsigned in_argc, char *in_argv[]) +{ + char **names = NULL; + const char *target_names = getenv (OFFLOAD_TARGET_NAMES_ENV); + if (!target_names) + return; + unsigned num_targets = parse_env_var (target_names, &names, NULL); + + const char *compiler_path = getenv ("COMPILER_PATH"); + if (!compiler_path) + goto out; + + /* Prepare an image for each target and save the name of the resultant object + file to the OFFLOAD_NAMES array. It is terminated by a NULL entry. */ + offload_names = XCNEWVEC (char *, num_targets + 1); + for (unsigned i = 0; i < num_targets; i++) + { + offload_names[i] = compile_offload_image (names[i], compiler_path, + in_argc, in_argv); + if (!offload_names[i]) + fatal_error ("problem with building target image for %s\n", names[i]); + } + + out: + 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 ("reading input file"); + if (len > 0) + { + fwrite (buffer, 1, len, d); + if (ferror (d) != 0) + fatal_error ("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 ("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 ("installation error, can't find crtoffloadbegin.o"); + + free_array_of_ptrs ((void **) paths, n_paths); +} + /* Execute gcc. ARGC is the number of arguments. ARGV contains the arguments. */ static void @@ -384,6 +607,8 @@ run_gcc (unsigned argc, char *argv[]) unsigned int decoded_options_count; struct obstack argv_obstack; int new_head_argc; + bool have_lto = false; + bool have_offload = false; /* Get the driver and options. */ collect_gcc = getenv ("COLLECT_GCC"); @@ -432,6 +657,9 @@ run_gcc (unsigned argc, char *argv[]) close (fd); continue; } + if (simple_object_find_section (sobj, OFFLOAD_SECTION_NAME_PREFIX ".opts", + &offset, &length, &errmsg, &err)) + have_offload = true; if (!simple_object_find_section (sobj, LTO_SECTION_NAME_PREFIX "." "opts", &offset, &length, &errmsg, &err)) { @@ -439,6 +667,7 @@ run_gcc (unsigned argc, char *argv[]) close (fd); continue; } + have_lto = true; lseek (fd, file_offset + offset, SEEK_SET); data = (char *)xmalloc (length); read (fd, data, length); @@ -633,6 +862,43 @@ run_gcc (unsigned argc, char *argv[]) /* Remember at which point we can scrub args to re-use the commons. */ new_head_argc = obstack_object_size (&argv_obstack) / sizeof (void *); + if (have_offload) + { + compile_images_for_offload_targets (argc, argv); + 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)) + { + 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 (lto_mode == LTO_MODE_LTO) { flto_out = make_temp_file (".lto.o"); @@ -859,6 +1125,10 @@ cont: obstack_free (&env_obstack, NULL); } + finish: + if (offloadend) + printf ("%s\n", offloadend); + obstack_free (&argv_obstack, NULL); } diff --git a/gcc/varasm.c b/gcc/varasm.c index abb743b..4ae9d58 100644 --- a/gcc/varasm.c +++ b/gcc/varasm.c @@ -6141,8 +6141,10 @@ default_elf_asm_named_section (const char *name, unsigned int flags, if (!(flags & SECTION_DEBUG)) *f++ = 'a'; +#if defined (HAVE_GAS_SECTION_EXCLUDE) && HAVE_GAS_SECTION_EXCLUDE == 1 if (flags & SECTION_EXCLUDE) *f++ = 'e'; +#endif if (flags & SECTION_WRITE) *f++ = 'w'; if (flags & SECTION_CODE) diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c index 910e23c..fb6555d 100644 --- a/lto-plugin/lto-plugin.c +++ b/lto-plugin/lto-plugin.c @@ -86,6 +86,8 @@ along with this program; see the file COPYING3. If not see #define LTO_SECTION_PREFIX ".gnu.lto_.symtab" #define LTO_SECTION_PREFIX_LEN (sizeof (LTO_SECTION_PREFIX) - 1) +#define OFFLOAD_SECTION ".gnu.offload_lto_.opts" +#define OFFLOAD_SECTION_LEN (sizeof (OFFLOAD_SECTION) - 1) /* The part of the symbol table the plugin has to keep track of. Note that we must keep SYMS until all_symbols_read is called to give the linker time to @@ -111,6 +113,7 @@ struct plugin_symtab struct plugin_objfile { int found; + int offload; simple_object_read *objfile; struct plugin_symtab *out; const struct ld_plugin_input_file *file; @@ -862,6 +865,21 @@ err: return 0; } +/* Find an offload section of an object file. */ + +static int +process_offload_section (void *data, const char *name, off_t offset, off_t len) +{ + if (!strncmp (name, OFFLOAD_SECTION, OFFLOAD_SECTION_LEN)) + { + struct plugin_objfile *obj = (struct plugin_objfile *) data; + obj->offload = 1; + return 0; + } + + return 1; +} + /* Callback used by gold to check if the plugin will claim FILE. Writes the result in CLAIMED. */ @@ -899,6 +917,7 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed) *claimed = 0; obj.file = file; obj.found = 0; + obj.offload = 0; obj.out = <o_file.symtab; errmsg = NULL; obj.objfile = simple_object_start_read (file->fd, file->offset, LTO_SEGMENT_NAME, @@ -920,7 +939,11 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed) goto err; } - if (obj.found == 0) + if (obj.objfile) + simple_object_find_sections (obj.objfile, process_offload_section, + &obj, &err); + + if (obj.found == 0 && obj.offload == 0) goto err; if (obj.found > 1) -- 1.7.1 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-15 14:27 ` Ilya Verbin @ 2014-10-15 14:46 ` Jakub Jelinek 2014-10-16 11:18 ` Ilya Verbin 0 siblings, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2014-10-15 14:46 UTC (permalink / raw) To: Ilya Verbin Cc: Bernd Schmidt, Cary Coutant, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On Wed, Oct 15, 2014 at 06:18:56PM +0400, Ilya Verbin wrote: > On 14 Oct 11:40, Jakub Jelinek wrote: > > My preference would be to add the | SECTION_EXCLUDE unconditionally, and > > instead guard the > > if (flags & SECTION_EXCLUDE) > > *f++ = 'e'; > > in varasm.c (default_elf_asm_named_section). The only other user of > > SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change > > ok with you? > > Done. But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac: > > gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],, > [.section foo1,"e" > .byte 0,0,0,0]) > > does not work properly. Maybe it works on "cygwin* | pe | mingw32* | interix*" > targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude > sections) it successfully assembles conftest.s into conftest.o (with warnings) > and HAVE_GAS_SECTION_EXCLUDE becomes defined. IMHO a version check is wrong (except when using in-tree gas). I'd suggest just to use [--fatal-warnings] as the 4th argument to gcc_GAS_CHECK_FEATURE feature, after all, that is what e.g. gcc_cv_as_shf_merge testing already uses. Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-15 14:46 ` Jakub Jelinek @ 2014-10-16 11:18 ` Ilya Verbin 2014-10-16 11:27 ` Jakub Jelinek 0 siblings, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-16 11:18 UTC (permalink / raw) To: Jakub Jelinek Cc: Bernd Schmidt, Cary Coutant, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On 15 Oct 16:38, Jakub Jelinek wrote: > > Done. But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac: > > > > gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],, > > [.section foo1,"e" > > .byte 0,0,0,0]) > > > > does not work properly. Maybe it works on "cygwin* | pe | mingw32* | interix*" > > targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude > > sections) it successfully assembles conftest.s into conftest.o (with warnings) > > and HAVE_GAS_SECTION_EXCLUDE becomes defined. > > IMHO a version check is wrong (except when using in-tree gas). > I'd suggest just to use [--fatal-warnings] as the 4th argument to > gcc_GAS_CHECK_FEATURE feature, after all, that is what e.g. > gcc_cv_as_shf_merge testing already uses. Fixed. Patch is updated and retested. Thanks, -- Ilya gcc/ * configure: Regenerate. * configure.ac: Move the test for section attribute specifier "e" in GAS out to all i[34567]86-*-* | x86_64-*-* targets and add --fatal-warnings. * gcc.c (spec_host_machine, accel_dir_suffix): New variables. (process_command): Tweak path construction for the possibility of being configured as an offload compiler. (driver::maybe_putenv_OFFLOAD_TARGETS): New function. (driver::main): Call maybe_putenv_OFFLOAD_TARGETS. (driver::set_up_specs): Tweak path construction for the possibility of being configured as an offload compiler. * langhooks.c (lhd_begin_section): Set SECTION_EXCLUDE flag. * lto-wrapper.c (OFFLOAD_TARGET_NAMES_ENV): Define. (offload_names, offloadbegin, offloadend): New static variables. (free_array_of_ptrs, parse_env_var, access_check, compile_offload_image) (compile_images_for_offload_targets, copy_file, find_offloadbeginend): New static functions. (run_gcc): Determine whether offload sections are present. If so, run compile_images_for_offload_targets and return the names of new generated objects to linker. If there are offload sections, but no LTO sections, then return the copies of input objects without link-time recompilation. * varasm.c (default_elf_asm_named_section): Guard SECTION_EXCLUDE with ifdef HAVE_GAS_SECTION_EXCLUDE. lto-plugin/ * lto-plugin.c (OFFLOAD_SECTION, OFFLOAD_SECTION_LEN): Define. (struct plugin_objfile): Add new field "offload". (process_offload_section): New static function. (claim_file_handler): Claim file if it contains offload sections. --- diff --git a/gcc/configure b/gcc/configure index ff1e398..4ef208c 100755 --- a/gcc/configure +++ b/gcc/configure @@ -24722,9 +24722,12 @@ $as_echo "$as_me: WARNING: LTO for $target requires binutils >= 2.20.1, but vers ;; esac fi - # Test if the assembler supports the section flag 'e' for specifying - # an excluded section. - { $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for .section with e" >&5 + ;; + esac + + # Test if the assembler supports the section flag 'e' for specifying + # an excluded section. + { $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for .section with e" >&5 $as_echo_n "checking assembler for .section with e... " >&6; } if test "${gcc_cv_as_section_has_e+set}" = set; then : $as_echo_n "(cached) " >&6 @@ -24737,7 +24740,7 @@ fi elif test x$gcc_cv_as != x; then $as_echo '.section foo1,"e" .byte 0,0,0,0' > conftest.s - if { ac_try='$gcc_cv_as $gcc_cv_as_flags -o conftest.o conftest.s >&5' + if { ac_try='$gcc_cv_as $gcc_cv_as_flags --fatal-warnings -o conftest.o conftest.s >&5' { { eval echo "\"\$as_me\":${as_lineno-$LINENO}: \"$ac_try\""; } >&5 (eval $ac_try) 2>&5 ac_status=$? @@ -24760,8 +24763,6 @@ cat >>confdefs.h <<_ACEOF #define HAVE_GAS_SECTION_EXCLUDE `if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi` _ACEOF - ;; - esac { $as_echo "$as_me:${as_lineno-$LINENO}: checking assembler for filds and fists mnemonics" >&5 $as_echo_n "checking assembler for filds and fists mnemonics... " >&6; } diff --git a/gcc/configure.ac b/gcc/configure.ac index 05a55f4..0f4bfc6 100644 --- a/gcc/configure.ac +++ b/gcc/configure.ac @@ -3837,18 +3837,19 @@ foo: nop ;; esac fi - # Test if the assembler supports the section flag 'e' for specifying - # an excluded section. - gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, - [2,22,51],, -[.section foo1,"e" -.byte 0,0,0,0]) - AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE, - [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`], - [Define if your assembler supports specifying the section flag e.]) ;; esac + # Test if the assembler supports the section flag 'e' for specifying + # an excluded section. + gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, + [2,22,51], [--fatal-warnings], +[.section foo1,"e" +.byte 0,0,0,0]) + AC_DEFINE_UNQUOTED(HAVE_GAS_SECTION_EXCLUDE, + [`if test $gcc_cv_as_section_has_e = yes; then echo 1; else echo 0; fi`], + [Define if your assembler supports specifying the section flag e.]) + gcc_GAS_CHECK_FEATURE([filds and fists mnemonics], gcc_cv_as_ix86_filds,,, [filds mem; fists mem],, diff --git a/gcc/gcc.c b/gcc/gcc.c index 71c76f8..4619fe7 100644 --- a/gcc/gcc.c +++ b/gcc/gcc.c @@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION; /* The target machine. */ static const char *spec_machine = DEFAULT_TARGET_MACHINE; +static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE; /* Nonzero if cross-compiling. When -b is used, the value comes from the `specs' file. */ @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2 relative to the driver. */ static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX; +/* A prefix to be used when this is an accelerator compiler. */ +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX; + /* Subdirectory to use for locating libraries. Set by set_multilib_dir based on the compilation options. */ @@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count, } gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); - tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, + tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, dir_separator_str, NULL); /* Look for tools relative to the location from which the driver is running, or, if that is not available, the configured prefix. */ tooldir_prefix = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix, - spec_machine, dir_separator_str, - spec_version, dir_separator_str, tooldir_prefix2, NULL); + spec_host_machine, dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL); free (tooldir_prefix2); add_prefix (&exec_prefixes, @@ -6742,6 +6746,7 @@ class driver void set_up_specs () const; void putenv_COLLECT_GCC (const char *argv0) const; void maybe_putenv_COLLECT_LTO_WRAPPER () const; + void maybe_putenv_OFFLOAD_TARGETS () const; void handle_unrecognized_options () const; int maybe_print_and_exit () const; bool prepare_infiles (); @@ -6784,6 +6789,7 @@ driver::main (int argc, char **argv) set_up_specs (); putenv_COLLECT_GCC (argv[0]); maybe_putenv_COLLECT_LTO_WRAPPER (); + maybe_putenv_OFFLOAD_TARGETS (); handle_unrecognized_options (); if (!maybe_print_and_exit ()) @@ -6953,6 +6959,7 @@ driver::build_multilib_strings () const void driver::set_up_specs () const { + const char *spec_machine_suffix; char *specs_file; size_t i; @@ -6976,8 +6983,8 @@ driver::set_up_specs () const /* Read specs from a file if there is one. */ - machine_suffix = concat (spec_machine, dir_separator_str, - spec_version, dir_separator_str, NULL); + machine_suffix = concat (spec_host_machine, dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, NULL); just_machine_suffix = concat (spec_machine, dir_separator_str, NULL); specs_file = find_a_file (&startfile_prefixes, "specs", R_OK, true); @@ -6987,13 +6994,18 @@ driver::set_up_specs () const else init_spec (); - /* We need to check standard_exec_prefix/just_machine_suffix/specs +#ifdef ACCEL_COMPILER + spec_machine_suffix = machine_suffix; +#else + spec_machine_suffix = just_machine_suffix; +#endif + + /* We need to check standard_exec_prefix/spec_machine_suffix/specs for any override of as, ld and libraries. */ specs_file = (char *) alloca (strlen (standard_exec_prefix) - + strlen (just_machine_suffix) + sizeof ("specs")); - + + strlen (spec_machine_suffix) + sizeof ("specs")); strcpy (specs_file, standard_exec_prefix); - strcat (specs_file, just_machine_suffix); + strcat (specs_file, spec_machine_suffix); strcat (specs_file, "specs"); if (access (specs_file, R_OK) == 0) read_specs (specs_file, true, false); @@ -7175,8 +7187,9 @@ driver::set_up_specs () const /* If we have a GCC_EXEC_PREFIX envvar, modify it for cpp's sake. */ if (gcc_exec_prefix) - gcc_exec_prefix = concat (gcc_exec_prefix, spec_machine, dir_separator_str, - spec_version, dir_separator_str, NULL); + gcc_exec_prefix = concat (gcc_exec_prefix, spec_host_machine, + dir_separator_str, spec_version, + accel_dir_suffix, dir_separator_str, NULL); /* Now we have the specs. Set the `valid' bits for switches that match anything in any spec. */ @@ -7227,6 +7240,21 @@ driver::maybe_putenv_COLLECT_LTO_WRAPPER () const } +/* Set up to remember the names of offload targets. */ + +void +driver::maybe_putenv_OFFLOAD_TARGETS () const +{ + if (strlen (OFFLOAD_TARGETS) > 0) + { + obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=", + sizeof ("OFFLOAD_TARGET_NAMES=") - 1); + obstack_grow (&collect_obstack, OFFLOAD_TARGETS, + strlen (OFFLOAD_TARGETS) + 1); + xputenv (XOBFINISH (&collect_obstack, char *)); + } +} + /* Reject switches that no pass was interested in. */ void diff --git a/gcc/langhooks.c b/gcc/langhooks.c index 7d4c294..4bdeaa0 100644 --- a/gcc/langhooks.c +++ b/gcc/langhooks.c @@ -660,7 +660,7 @@ lhd_begin_section (const char *name) saved_section = text_section; /* Create a new section and switch to it. */ - section = get_section (name, SECTION_DEBUG, NULL); + section = get_section (name, SECTION_DEBUG | SECTION_EXCLUDE, NULL); switch_to_section (section); } diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c index 8033b15..cbda36b 100644 --- a/gcc/lto-wrapper.c +++ b/gcc/lto-wrapper.c @@ -49,6 +49,10 @@ along with GCC; see the file COPYING3. If not see #include "lto-section-names.h" #include "collect-utils.h" +/* Environment variable, used for passing the names of offload targets from GCC + driver to lto-wrapper. */ +#define OFFLOAD_TARGET_NAMES_ENV "OFFLOAD_TARGET_NAMES" + enum lto_mode_d { LTO_MODE_NONE, /* Not doing LTO. */ LTO_MODE_LTO, /* Normal LTO. */ @@ -63,6 +67,8 @@ static char *flto_out; static unsigned int nr; static char **input_names; static char **output_names; +static char **offload_names; +static const char *offloadbegin, *offloadend; static char *makefile; const char tool_name[] = "lto-wrapper"; @@ -364,6 +370,223 @@ merge_and_complain (struct cl_decoded_option **decoded_options, } } +/* Auxiliary function that frees elements of PTR and PTR itself. + N is number of elements to be freed. If PTR is NULL, nothing is freed. + If an element is NULL, subsequent elements are not freed. */ + +static void ** +free_array_of_ptrs (void **ptr, unsigned n) +{ + if (!ptr) + return NULL; + for (unsigned i = 0; i < n; i++) + { + if (!ptr[i]) + break; + free (ptr[i]); + } + free (ptr); + return NULL; +} + +/* Parse STR, saving found tokens into PVALUES and return their number. + Tokens are assumed to be delimited by ':'. If APPEND is non-null, + append it to every token we find. */ + +static unsigned +parse_env_var (const char *str, char ***pvalues, const char *append) +{ + const char *curval, *nextval; + char **values; + unsigned num = 1, i; + + curval = strchr (str, ':'); + while (curval) + { + num++; + curval = strchr (curval + 1, ':'); + } + + values = (char**) xmalloc (num * sizeof (char*)); + curval = str; + nextval = strchrnul (curval, ':'); + + int append_len = append ? strlen (append) : 0; + for (i = 0; i < num; i++) + { + int l = nextval - curval; + values[i] = (char*) xmalloc (l + 1 + append_len); + memcpy (values[i], curval, l); + values[i][l] = 0; + if (append) + strcat (values[i], append); + curval = nextval + 1; + nextval = strchrnul (curval, ':'); + } + *pvalues = values; + return num; +} + +/* Check whether NAME can be accessed in MODE. This is like access, + except that it never considers directories to be executable. */ + +static int +access_check (const char *name, int mode) +{ + if (mode == X_OK) + { + struct stat st; + + if (stat (name, &st) < 0 + || S_ISDIR (st.st_mode)) + return -1; + } + + return access (name, mode); +} + +/* Prepare a target image for offload TARGET, using mkoffload tool from + COMPILER_PATH. Return the name of the resultant object file. */ + +static char * +compile_offload_image (const char *target, const char *compiler_path, + unsigned in_argc, char *in_argv[]) +{ + char *filename = NULL; + char **argv; + char *suffix + = XALLOCAVEC (char, sizeof ("/accel//mkoffload") + strlen (target)); + strcpy (suffix, "/accel/"); + strcat (suffix, target); + strcat (suffix, "/mkoffload"); + + char **paths = NULL; + unsigned n_paths = parse_env_var (compiler_path, &paths, suffix); + + const char *compiler = NULL; + for (unsigned i = 0; i < n_paths; i++) + if (access_check (paths[i], X_OK) == 0) + { + compiler = paths[i]; + break; + } + + if (compiler) + { + /* Generate temporary output file name. */ + filename = make_temp_file (".target.o"); + + struct obstack argv_obstack; + obstack_init (&argv_obstack); + obstack_ptr_grow (&argv_obstack, compiler); + obstack_ptr_grow (&argv_obstack, "-o"); + obstack_ptr_grow (&argv_obstack, filename); + + for (unsigned i = 1; i < in_argc; i++) + obstack_ptr_grow (&argv_obstack, in_argv[i]); + obstack_ptr_grow (&argv_obstack, NULL); + + argv = XOBFINISH (&argv_obstack, char **); + fork_execute (argv[0], argv, true); + obstack_free (&argv_obstack, NULL); + } + + free_array_of_ptrs ((void **) paths, n_paths); + return filename; +} + + +/* The main routine dealing with offloading. + The routine builds a target image for each offload target. IN_ARGC and + IN_ARGV specify options and input object files. As all of them could contain + target sections, we pass them all to target compilers. */ + +static void +compile_images_for_offload_targets (unsigned in_argc, char *in_argv[]) +{ + char **names = NULL; + const char *target_names = getenv (OFFLOAD_TARGET_NAMES_ENV); + if (!target_names) + return; + unsigned num_targets = parse_env_var (target_names, &names, NULL); + + const char *compiler_path = getenv ("COMPILER_PATH"); + if (!compiler_path) + goto out; + + /* Prepare an image for each target and save the name of the resultant object + file to the OFFLOAD_NAMES array. It is terminated by a NULL entry. */ + offload_names = XCNEWVEC (char *, num_targets + 1); + for (unsigned i = 0; i < num_targets; i++) + { + offload_names[i] = compile_offload_image (names[i], compiler_path, + in_argc, in_argv); + if (!offload_names[i]) + fatal_error ("problem with building target image for %s\n", names[i]); + } + + out: + 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 ("reading input file"); + if (len > 0) + { + fwrite (buffer, 1, len, d); + if (ferror (d) != 0) + fatal_error ("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 ("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 ("installation error, can't find crtoffloadbegin.o"); + + free_array_of_ptrs ((void **) paths, n_paths); +} + /* Execute gcc. ARGC is the number of arguments. ARGV contains the arguments. */ static void @@ -384,6 +607,8 @@ run_gcc (unsigned argc, char *argv[]) unsigned int decoded_options_count; struct obstack argv_obstack; int new_head_argc; + bool have_lto = false; + bool have_offload = false; /* Get the driver and options. */ collect_gcc = getenv ("COLLECT_GCC"); @@ -432,6 +657,9 @@ run_gcc (unsigned argc, char *argv[]) close (fd); continue; } + if (simple_object_find_section (sobj, OFFLOAD_SECTION_NAME_PREFIX ".opts", + &offset, &length, &errmsg, &err)) + have_offload = true; if (!simple_object_find_section (sobj, LTO_SECTION_NAME_PREFIX "." "opts", &offset, &length, &errmsg, &err)) { @@ -439,6 +667,7 @@ run_gcc (unsigned argc, char *argv[]) close (fd); continue; } + have_lto = true; lseek (fd, file_offset + offset, SEEK_SET); data = (char *)xmalloc (length); read (fd, data, length); @@ -633,6 +862,43 @@ run_gcc (unsigned argc, char *argv[]) /* Remember at which point we can scrub args to re-use the commons. */ new_head_argc = obstack_object_size (&argv_obstack) / sizeof (void *); + if (have_offload) + { + compile_images_for_offload_targets (argc, argv); + 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)) + { + 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 (lto_mode == LTO_MODE_LTO) { flto_out = make_temp_file (".lto.o"); @@ -859,6 +1125,10 @@ cont: obstack_free (&env_obstack, NULL); } + finish: + if (offloadend) + printf ("%s\n", offloadend); + obstack_free (&argv_obstack, NULL); } diff --git a/gcc/varasm.c b/gcc/varasm.c index abb743b..4ae9d58 100644 --- a/gcc/varasm.c +++ b/gcc/varasm.c @@ -6141,8 +6141,10 @@ default_elf_asm_named_section (const char *name, unsigned int flags, if (!(flags & SECTION_DEBUG)) *f++ = 'a'; +#if defined (HAVE_GAS_SECTION_EXCLUDE) && HAVE_GAS_SECTION_EXCLUDE == 1 if (flags & SECTION_EXCLUDE) *f++ = 'e'; +#endif if (flags & SECTION_WRITE) *f++ = 'w'; if (flags & SECTION_CODE) diff --git a/lto-plugin/lto-plugin.c b/lto-plugin/lto-plugin.c index 910e23c..fb6555d 100644 --- a/lto-plugin/lto-plugin.c +++ b/lto-plugin/lto-plugin.c @@ -86,6 +86,8 @@ along with this program; see the file COPYING3. If not see #define LTO_SECTION_PREFIX ".gnu.lto_.symtab" #define LTO_SECTION_PREFIX_LEN (sizeof (LTO_SECTION_PREFIX) - 1) +#define OFFLOAD_SECTION ".gnu.offload_lto_.opts" +#define OFFLOAD_SECTION_LEN (sizeof (OFFLOAD_SECTION) - 1) /* The part of the symbol table the plugin has to keep track of. Note that we must keep SYMS until all_symbols_read is called to give the linker time to @@ -111,6 +113,7 @@ struct plugin_symtab struct plugin_objfile { int found; + int offload; simple_object_read *objfile; struct plugin_symtab *out; const struct ld_plugin_input_file *file; @@ -862,6 +865,21 @@ err: return 0; } +/* Find an offload section of an object file. */ + +static int +process_offload_section (void *data, const char *name, off_t offset, off_t len) +{ + if (!strncmp (name, OFFLOAD_SECTION, OFFLOAD_SECTION_LEN)) + { + struct plugin_objfile *obj = (struct plugin_objfile *) data; + obj->offload = 1; + return 0; + } + + return 1; +} + /* Callback used by gold to check if the plugin will claim FILE. Writes the result in CLAIMED. */ @@ -899,6 +917,7 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed) *claimed = 0; obj.file = file; obj.found = 0; + obj.offload = 0; obj.out = <o_file.symtab; errmsg = NULL; obj.objfile = simple_object_start_read (file->fd, file->offset, LTO_SEGMENT_NAME, @@ -920,7 +939,11 @@ claim_file_handler (const struct ld_plugin_input_file *file, int *claimed) goto err; } - if (obj.found == 0) + if (obj.objfile) + simple_object_find_sections (obj.objfile, process_offload_section, + &obj, &err); + + if (obj.found == 0 && obj.offload == 0) goto err; if (obj.found > 1) -- 1.7.1 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-16 11:18 ` Ilya Verbin @ 2014-10-16 11:27 ` Jakub Jelinek 2014-10-29 10:28 ` Kirill Yukhin 2014-11-06 13:00 ` Ilya Verbin 0 siblings, 2 replies; 24+ messages in thread From: Jakub Jelinek @ 2014-10-16 11:27 UTC (permalink / raw) To: Ilya Verbin Cc: Bernd Schmidt, Cary Coutant, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy On Thu, Oct 16, 2014 at 03:17:36PM +0400, Ilya Verbin wrote: > On 15 Oct 16:38, Jakub Jelinek wrote: > > > Done. But it turned out that the gcc_GAS_CHECK_FEATURE from gcc/configure.ac: > > > > > > gcc_GAS_CHECK_FEATURE([.section with e], gcc_cv_as_section_has_e, [2,22,51],, > > > [.section foo1,"e" > > > .byte 0,0,0,0]) > > > > > > does not work properly. Maybe it works on "cygwin* | pe | mingw32* | interix*" > > > targets, but on linux with GNU as version 2.20.51 (which doesn't support exclude > > > sections) it successfully assembles conftest.s into conftest.o (with warnings) > > > and HAVE_GAS_SECTION_EXCLUDE becomes defined. > > > > IMHO a version check is wrong (except when using in-tree gas). > > I'd suggest just to use [--fatal-warnings] as the 4th argument to > > gcc_GAS_CHECK_FEATURE feature, after all, that is what e.g. > > gcc_cv_as_shf_merge testing already uses. > > Fixed. Patch is updated and retested. Can you please extract the configure{,.ac}, langhooks.c and varasm.c bits into a separate patch? That is preapproved for trunk right now, that isn't dependent on anything else. The rest LGTM, but please run it through LTO review (Richard/Honza) too. Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-16 11:27 ` Jakub Jelinek @ 2014-10-29 10:28 ` Kirill Yukhin 2014-11-06 13:00 ` Ilya Verbin 1 sibling, 0 replies; 24+ messages in thread From: Kirill Yukhin @ 2014-10-29 10:28 UTC (permalink / raw) To: Richard Biener, Jan Hubicka Cc: Ilya Verbin, Bernd Schmidt, Cary Coutant, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Andrey Turetskiy, Jakub Jelinek Hello Richard, Jan, On 16 Oct 13:22, Jakub Jelinek wrote: > On Thu, Oct 16, 2014 at 03:17:36PM +0400, Ilya Verbin wrote: > The rest LGTM, but please run it through LTO review (Richard/Honza) too. Ping? -- Thanks, k > > Jakub ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-16 11:27 ` Jakub Jelinek 2014-10-29 10:28 ` Kirill Yukhin @ 2014-11-06 13:00 ` Ilya Verbin 2014-11-12 9:47 ` Richard Biener 1 sibling, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-11-06 13:00 UTC (permalink / raw) To: Richard Biener Cc: Jan Hubicka, Jakub Jelinek, Jeff Law, Bernd Schmidt, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy, gcc-patches On 16 Oct 13:22, Jakub Jelinek wrote: > Can you please extract the configure{,.ac}, langhooks.c and varasm.c > bits into a separate patch? That is preapproved for trunk right now, that > isn't dependent on anything else. Done. > The rest LGTM, but please run it through LTO review (Richard/Honza) too. Richard, is this patch OK for trunk? Thanks, -- Ilya ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-11-06 13:00 ` Ilya Verbin @ 2014-11-12 9:47 ` Richard Biener 0 siblings, 0 replies; 24+ messages in thread From: Richard Biener @ 2014-11-12 9:47 UTC (permalink / raw) To: Ilya Verbin Cc: Jan Hubicka, Jakub Jelinek, Jeff Law, Bernd Schmidt, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy, gcc-patches On Thu, 6 Nov 2014, Ilya Verbin wrote: > On 16 Oct 13:22, Jakub Jelinek wrote: > > Can you please extract the configure{,.ac}, langhooks.c and varasm.c > > bits into a separate patch? That is preapproved for trunk right now, that > > isn't dependent on anything else. > > Done. > > > The rest LGTM, but please run it through LTO review (Richard/Honza) too. > > Richard, is this patch OK for trunk? Ok. Thanks, Richard. ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-14 9:41 ` Jakub Jelinek 2014-10-15 14:27 ` Ilya Verbin @ 2014-10-15 17:03 ` Cary Coutant 1 sibling, 0 replies; 24+ messages in thread From: Cary Coutant @ 2014-10-15 17:03 UTC (permalink / raw) To: Jakub Jelinek Cc: Ilya Verbin, Richard Biener, Jan Hubicka, gcc-patches, Thomas Schwinge, Kirill Yukhin, Andrey Turetskiy > My preference would be to add the | SECTION_EXCLUDE unconditionally, and > instead guard the > if (flags & SECTION_EXCLUDE) > *f++ = 'e'; > in varasm.c (default_elf_asm_named_section). The only other user of > SECTION_EXCLUDE seems to be -gsplit-dwarf right now, Cary, is such a change > ok with you? Yes, that sounds fine. > If you have new gas and old linker, I'd expect it would just ignore > SHF_EXCLUDE. Agreed. -cary ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-02 15:15 [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Ilya Verbin 2014-10-08 10:27 ` Jakub Jelinek @ 2015-05-12 16:32 ` Thomas Schwinge 2015-05-12 17:19 ` Bernd Schmidt 2016-02-19 19:42 ` [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Thomas Schwinge 2 siblings, 1 reply; 24+ messages in thread From: Thomas Schwinge @ 2015-05-12 16:32 UTC (permalink / raw) To: Bernd Schmidt Cc: Kirill Yukhin, Andrey Turetskiy, Ilya Verbin, Jakub Jelinek, Jan Hubicka, Richard Biener, gcc-patches [-- Attachment #1: Type: text/plain, Size: 4745 bytes --] Hi Bernd! One question to you, as the one who has originally written this code (<https://gcc.gnu.org/ml/gcc-patches/2014-03/msg01079.html>): On Thu, 2 Oct 2014 19:14:57 +0400, Ilya Verbin <iverbin@gmail.com> wrote: > With this patch lto-wrapper performs invocation of mkoffload tool for each > offload target. [...] (This has been committed long ago.) > 2014-10-02 Ilya Verbin <ilya.verbin@intel.com> > Bernd Schmidt <bernds@codesourcery.com> > Andrey Turetskiy <andrey.turetskiy@intel.com> > Michael Zolotukhin <michael.v.zolotukhin@intel.com> > > gcc/ > * gcc.c (spec_host_machine, accel_dir_suffix): New variables. > (process_command): Tweak path construction for the possibility > of being configured as an offload compiler. > (main): Likewise. Look up specs in just_machine_suffix only if not > ACCEL_COMPILER. Construct OFFLOAD_TARGET_NAMES environment variable if > we have OFFLOAD_TARGETS. > * [...] > --- a/gcc/gcc.c > +++ b/gcc/gcc.c > @@ -157,6 +157,7 @@ static const char *const spec_version = DEFAULT_TARGET_VERSION; > /* The target machine. */ > > static const char *spec_machine = DEFAULT_TARGET_MACHINE; > +static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE; > > /* Nonzero if cross-compiling. > When -b is used, the value comes from the `specs' file. */ > @@ -1296,6 +1297,9 @@ static const char *const standard_startfile_prefix_2 > relative to the driver. */ > static const char *const tooldir_base_prefix = TOOLDIR_BASE_PREFIX; > > +/* A prefix to be used when this is an accelerator compiler. */ > +static const char *const accel_dir_suffix = ACCEL_DIR_SUFFIX; > + > /* Subdirectory to use for locating libraries. Set by > set_multilib_dir based on the compilation options. */ > > @@ -4122,15 +4126,15 @@ process_command (unsigned int decoded_options_count, > } > > gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); > - tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, > + tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, > dir_separator_str, NULL); I've noticed that internally we have a change (r447328) to »Use spec_machine rather than spec_host_machine to build tooldir_prefix2«, thus you reverted this last change, describing this as a »Driver paths bugfix«, which »corrects an issue that causes an offload gcc driver to find the wrong assembler«. In addition to that change/reversion, one of your more recent patche submissions, <https://gcc.gnu.org/ml/gcc-patches/2014-11/msg02655.html>, contained the (unrelated) change to also change back to the original spec_machine the spec_host_machine usage in the following code: > > /* Look for tools relative to the location from which the driver is > running, or, if that is not available, the configured prefix. */ > tooldir_prefix > = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix, > - spec_machine, dir_separator_str, > - spec_version, dir_separator_str, tooldir_prefix2, NULL); > + spec_host_machine, dir_separator_str, spec_version, > + accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL); > free (tooldir_prefix2); > > add_prefix (&exec_prefixes, Which of the changes should we be making on trunk? None at all, or just revert the change for tooldir_prefix2 (patch variant 1), or also for tooldir_prefix (patch variant 2), or something else? Patch variant 1: @@ -4266,7 +4266,7 @@ process_command (unsigned int decoded_op } gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); - tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, + tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, dir_separator_str, NULL); /* Look for tools relative to the location from which the driver is Patch variant 2: @@ -4238,14 +4238,14 @@ process_command (unsigned int decoded_options_count, } gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); - tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, + tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, dir_separator_str, NULL); /* Look for tools relative to the location from which the driver is running, or, if that is not available, the configured prefix. */ tooldir_prefix = concat (gcc_exec_prefix ? gcc_exec_prefix : standard_exec_prefix, - spec_host_machine, dir_separator_str, spec_version, + spec_machine, dir_separator_str, spec_version, accel_dir_suffix, dir_separator_str, tooldir_prefix2, NULL); free (tooldir_prefix2); Grüße, Thomas [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2015-05-12 16:32 ` Thomas Schwinge @ 2015-05-12 17:19 ` Bernd Schmidt 2015-10-02 21:28 ` Help the offload gcc driver find the right assembler (was: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper) Thomas Schwinge 0 siblings, 1 reply; 24+ messages in thread From: Bernd Schmidt @ 2015-05-12 17:19 UTC (permalink / raw) To: Thomas Schwinge Cc: Kirill Yukhin, Andrey Turetskiy, Ilya Verbin, Jakub Jelinek, Jan Hubicka, Richard Biener, gcc-patches On 05/12/2015 06:27 PM, Thomas Schwinge wrote: > Patch variant 1: > > @@ -4266,7 +4266,7 @@ process_command (unsigned int decoded_op > } > > gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); > - tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, > + tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, > dir_separator_str, NULL); > > /* Look for tools relative to the location from which the driver is This one, definitely. The other one I can't remember. Are there any testsuite results for either variant? Bernd ^ permalink raw reply [flat|nested] 24+ messages in thread
* Help the offload gcc driver find the right assembler (was: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper) 2015-05-12 17:19 ` Bernd Schmidt @ 2015-10-02 21:28 ` Thomas Schwinge 0 siblings, 0 replies; 24+ messages in thread From: Thomas Schwinge @ 2015-10-02 21:28 UTC (permalink / raw) To: Bernd Schmidt, gcc-patches Cc: Kirill Yukhin, Andrey Turetskiy, Ilya Verbin, Jakub Jelinek, Jan Hubicka, Richard Biener, nathan, James Norris [-- Attachment #1: Type: text/plain, Size: 2281 bytes --] Hi! On Tue, 12 May 2015 19:05:17 +0200, Bernd Schmidt <bernds@codesourcery.com> wrote: > > @@ -4266,7 +4266,7 @@ process_command (unsigned int decoded_op > > } > > > > gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); > > - tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, > > + tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, > > dir_separator_str, NULL); > > > > /* Look for tools relative to the location from which the driver is > > This one, definitely. Indeed; now committed in r228429. (Jim, did this missing patch cause the thousands of FAILs you've seen? The problem was that nvptx offloading found the wrong assembler, and we all know what happens when you try to process PTX assembly with an x86/PowerPC assembler...) commit f2fd5997e0856128c0609cdc6aa4cf9867f94c41 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Fri Oct 2 21:24:50 2015 +0000 Help the offload gcc driver find the right assembler gcc/ * gcc.c (process_command): Use spec_machine rather than spec_host_machine to build tooldir_prefix2. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@228429 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog | 5 +++++ gcc/gcc.c | 2 +- 2 files changed, 6 insertions(+), 1 deletion(-) diff --git gcc/ChangeLog gcc/ChangeLog index d1235bd..7fe0196 100644 --- gcc/ChangeLog +++ gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-10-02 Bernd Schmidt <bernds@codesourcery.com> + + * gcc.c (process_command): Use spec_machine rather than + spec_host_machine to build tooldir_prefix2. + 2015-10-02 Thomas Schwinge <thomas@codesourcery.com> Bernd Schmidt <bernds@codesourcery.com> diff --git gcc/gcc.c gcc/gcc.c index db90796..7f5a36e 100644 --- gcc/gcc.c +++ gcc/gcc.c @@ -4472,7 +4472,7 @@ process_command (unsigned int decoded_options_count, } gcc_assert (!IS_ABSOLUTE_PATH (tooldir_base_prefix)); - tooldir_prefix2 = concat (tooldir_base_prefix, spec_host_machine, + tooldir_prefix2 = concat (tooldir_base_prefix, spec_machine, dir_separator_str, NULL); /* Look for tools relative to the location from which the driver is Grüße, Thomas [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2014-10-02 15:15 [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Ilya Verbin 2014-10-08 10:27 ` Jakub Jelinek 2015-05-12 16:32 ` Thomas Schwinge @ 2016-02-19 19:42 ` Thomas Schwinge 2016-02-19 19:51 ` Ilya Verbin 2 siblings, 1 reply; 24+ messages in thread From: Thomas Schwinge @ 2016-02-19 19:42 UTC (permalink / raw) To: Ilya Verbin, Jakub Jelinek Cc: Bernd Schmidt, Kirill Yukhin, Andrey Turetskiy, Jan Hubicka, Richard Biener, gcc-patches, Cesar Philippidis [-- Attachment #1: Type: text/plain, Size: 4714 bytes --] Hi! On Thu, 2 Oct 2014 19:14:57 +0400, Ilya Verbin <iverbin@gmail.com> wrote: > With this patch lto-wrapper performs invocation of mkoffload tool for each > offload target. This tool [...] > will compile IR from .gnu.offload_lto_* sections into offload > target code and embed the resultant code (offload image) into the new host's > object file. Consider the following scenario: $ cat < CSTS-214-acc.c int acc (void) { int a; #pragma acc parallel num_gangs (1) copyout (a) a = 100; return a; } $ cat < CSTS-214-test.c extern int acc (void); int main (void) { if (acc () != 100) __builtin_abort (); return 0; } Compile these two files as follows: $ [GCC] -fopenacc -c CSTS-214-acc.c $ x86_64-linux-gnu-ar -cr CSTS-214-acc.a CSTS-214-acc.o $ [GCC] -fopenacc CSTS-214-test.c CSTS-214-acc.a The last step will fail -- with incomprehensible diagnostics, ;-) as so often when offloading fails... Here's what's going on: the LTO/offloading machinery correctly identifies that it needs to process the CSTS-214-acc.c:acc function, present in the CSTS-214-acc.a archive file at a certain offset, and it "encodes" that as follows: CSTS-214-acc.a@0x9e (see lto-plugin/lto-plugin.c:claim_file_handler, the "file->offset != 0" code right at the beginning). This makes its way down through here: > --- a/gcc/lto-wrapper.c > +++ b/gcc/lto-wrapper.c > +/* Copy a file from SRC to DEST. */ > + > +static void > +copy_file (const char *dest, const char *src) > +{ > + [...] > +} > @@ -624,6 +852,54 @@ run_gcc (unsigned argc, char *argv[]) > + /* 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 ([...]) > + { > + 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 (lto_mode == LTO_MODE_LTO) > { > flto_out = make_temp_file (".lto.o"); > @@ -850,6 +1126,10 @@ cont: > obstack_free (&env_obstack, NULL); > } > > + finish: > + if (offloadend) > + printf ("%s\n", offloadend); > + > obstack_free (&argv_obstack, NULL); > } When we hit this, for argv "CSTS-214-acc.a@0x9e", the copy_file call will fail -- there is no "CSTS-214-acc.a@0x9e" file to copy. If we strip off the "@0x[...]" suffix (but still printf the filename including the suffix), then things work. I copied that bit of code from earlier in this function, where the same archive offset handling needs to be done. Probably that code should be refactored a bit. Also, I wonder if the "ext == NULL" case can really happen, and needs to be handled as done in the code cited above, or if that can be simplified? (Not yet tested that.) Will something like the following be OK to fix this issue, or is that something "that should not happen", should be fixed differently? --- gcc/lto-wrapper.c +++ gcc/lto-wrapper.c @@ -1161,15 +1161,31 @@ run_gcc (unsigned argc, char *argv[]) && strncmp (argv[i], "-flinker-output=", sizeof ("-flinker-output=") - 1) != 0) { + char *p; + off_t file_offset = 0; + long loffset; + int consumed; + char *filename = argv[i]; + + if ((p = strrchr (argv[i], '@')) + && p != argv[i] + && sscanf (p, "@%li%n", &loffset, &consumed) >= 1 + && strlen (p) == (unsigned int) consumed) + { + filename = XNEWVEC (char, p - argv[i] + 1); + memcpy (filename, argv[i], p - argv[i]); + filename[p - argv[i]] = '\0'; + file_offset = (off_t) loffset; + } + char *out_file; - /* Can be ".o" or ".so". */ - char *ext = strrchr (argv[i], '.'); + char *ext = strrchr (filename, '.'); 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]); + copy_file (out_file, filename); printf ("%s\n", out_file); } goto finish; Grüße Thomas [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper 2016-02-19 19:42 ` [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Thomas Schwinge @ 2016-02-19 19:51 ` Ilya Verbin 0 siblings, 0 replies; 24+ messages in thread From: Ilya Verbin @ 2016-02-19 19:51 UTC (permalink / raw) To: Thomas Schwinge Cc: Jakub Jelinek, Bernd Schmidt, Kirill Yukhin, Richard Biener, gcc-patches On Fri, Feb 19, 2016 at 20:41:58 +0100, Thomas Schwinge wrote: > Hi! > > On Thu, 2 Oct 2014 19:14:57 +0400, Ilya Verbin <iverbin@gmail.com> wrote: > > With this patch lto-wrapper performs invocation of mkoffload tool for each > > offload target. This tool [...] > > will compile IR from .gnu.offload_lto_* sections into offload > > target code and embed the resultant code (offload image) into the new host's > > object file. > > Consider the following scenario: > > $ cat < CSTS-214-acc.c > int acc (void) > { > int a; > > #pragma acc parallel num_gangs (1) copyout (a) > a = 100; > > return a; > } > $ cat < CSTS-214-test.c > extern int acc (void); > > int main (void) > { > if (acc () != 100) > __builtin_abort (); > > return 0; > } > > Compile these two files as follows: > > $ [GCC] -fopenacc -c CSTS-214-acc.c > $ x86_64-linux-gnu-ar -cr CSTS-214-acc.a CSTS-214-acc.o > $ [GCC] -fopenacc CSTS-214-test.c CSTS-214-acc.a > > The last step will fail -- with incomprehensible diagnostics, ;-) as so > often when offloading fails... Here's what's going on: the > LTO/offloading machinery correctly identifies that it needs to process > the CSTS-214-acc.c:acc function, present in the CSTS-214-acc.a archive > file at a certain offset, and it "encodes" that as follows: > CSTS-214-acc.a@0x9e (see lto-plugin/lto-plugin.c:claim_file_handler, the > "file->offset != 0" code right at the beginning). This makes its way > down through here: > > > --- a/gcc/lto-wrapper.c > > +++ b/gcc/lto-wrapper.c > > > +/* Copy a file from SRC to DEST. */ > > + > > +static void > > +copy_file (const char *dest, const char *src) > > +{ > > + [...] > > +} > > > @@ -624,6 +852,54 @@ run_gcc (unsigned argc, char *argv[]) > > > + /* 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 ([...]) > > + { > > + 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 (lto_mode == LTO_MODE_LTO) > > { > > flto_out = make_temp_file (".lto.o"); > > @@ -850,6 +1126,10 @@ cont: > > obstack_free (&env_obstack, NULL); > > } > > > > + finish: > > + if (offloadend) > > + printf ("%s\n", offloadend); > > + > > obstack_free (&argv_obstack, NULL); > > } > > When we hit this, for argv "CSTS-214-acc.a@0x9e", the copy_file call will > fail -- there is no "CSTS-214-acc.a@0x9e" file to copy. If we strip off > the "@0x[...]" suffix (but still printf the filename including the > suffix), then things work. I copied that bit of code from earlier in > this function, where the same archive offset handling needs to be done. > Probably that code should be refactored a bit. > > Also, I wonder if the "ext == NULL" case can really happen, and needs to > be handled as done in the code cited above, or if that can be simplified? > (Not yet tested that.) > > Will something like the following be OK to fix this issue, or is that > something "that should not happen", should be fixed differently? > > --- gcc/lto-wrapper.c > +++ gcc/lto-wrapper.c > @@ -1161,15 +1161,31 @@ run_gcc (unsigned argc, char *argv[]) > && strncmp (argv[i], "-flinker-output=", > sizeof ("-flinker-output=") - 1) != 0) > { > + char *p; > + off_t file_offset = 0; > + long loffset; > + int consumed; > + char *filename = argv[i]; > + > + if ((p = strrchr (argv[i], '@')) > + && p != argv[i] > + && sscanf (p, "@%li%n", &loffset, &consumed) >= 1 > + && strlen (p) == (unsigned int) consumed) > + { > + filename = XNEWVEC (char, p - argv[i] + 1); > + memcpy (filename, argv[i], p - argv[i]); > + filename[p - argv[i]] = '\0'; > + file_offset = (off_t) loffset; > + } > + > char *out_file; > - /* Can be ".o" or ".so". */ > - char *ext = strrchr (argv[i], '.'); > + char *ext = strrchr (filename, '.'); > 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]); > + copy_file (out_file, filename); > printf ("%s\n", out_file); > } > goto finish; Yes, this part of lto-wrapper is awfully. This patch completely reworks it: https://gcc.gnu.org/ml/gcc-patches/2016-02/msg00709.html It's not yet fully ready, I'm going to send the final patch for review tomorrow. -- Ilya ^ permalink raw reply [flat|nested] 24+ messages in thread
end of thread, other threads:[~2016-02-19 19:51 UTC | newest] Thread overview: 24+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2014-10-02 15:15 [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Ilya Verbin 2014-10-08 10:27 ` Jakub Jelinek 2014-10-09 12:09 ` Ilya Verbin 2014-10-09 12:13 ` Bernd Schmidt 2014-10-09 20:27 ` Ilya Verbin 2014-10-10 7:13 ` Jakub Jelinek 2014-10-10 16:52 ` Cary Coutant 2014-10-10 17:01 ` Jakub Jelinek 2014-10-10 17:10 ` Cary Coutant 2014-10-13 22:47 ` Ilya Verbin 2014-10-14 9:41 ` Jakub Jelinek 2014-10-15 14:27 ` Ilya Verbin 2014-10-15 14:46 ` Jakub Jelinek 2014-10-16 11:18 ` Ilya Verbin 2014-10-16 11:27 ` Jakub Jelinek 2014-10-29 10:28 ` Kirill Yukhin 2014-11-06 13:00 ` Ilya Verbin 2014-11-12 9:47 ` Richard Biener 2014-10-15 17:03 ` Cary Coutant 2015-05-12 16:32 ` Thomas Schwinge 2015-05-12 17:19 ` Bernd Schmidt 2015-10-02 21:28 ` Help the offload gcc driver find the right assembler (was: [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper) Thomas Schwinge 2016-02-19 19:42 ` [PATCH 4/n] OpenMP 4.0 offloading infrastructure: lto-wrapper Thomas Schwinge 2016-02-19 19:51 ` Ilya Verbin
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).