[RFC,Offloading] Fix PR68463

Hi!

Here is my attempt to fix https://gcc.gnu.org/bugzilla/show_bug.cgi?id=68463

This patch does 2 things:

I) lto-plugin doesn't claim files which contain offload sections, but don't
contain LTO sections.  Instead, it writes names of files with offloading to the
temporary file and passes it to lto-wrapper as -foffload-objects=/tmp/cc...
The order of these files in the list is very important, because ld will link
host objects (and therefore host tables) in the following order:
  1. Non-LTO files before the first claimed LTO file;
  2. LTO files, after WPA-partitioning-recompilation;
  3. Non-LTO files after the first claimed LTO file.
To get the correct matching between host and target tables, the offload objects
need to be reordered correspondingly before passing to the target compiler.

II) The __offload_func_table, __offload_funcs_end, __offload_var_table,
__offload_vars_end are now provided by the linker script, instead of
crtoffload{begin,end}.o, this allows to surround all offload objects, even
those that are not claimed by lto-plugin.
Unfortunately it works only with ld, but doen't work with gold, because
https://sourceware.org/bugzilla/show_bug.cgi?id=15373
Any thoughts how to enable this linker script for gold?

I used the following testcase:
$ cat main.c
void foo1 ();
void foo2 ();
void foo3 ();
void foo4 ();

int main ()
{
  foo1 ();
  foo2 ();
  foo3 ();
  foo4 ();
  return 0;
}

$ cat test.c
#include <stdio.h>
#include <omp.h>
#define MAKE_FN_NAME(x) foo ## x
#define FN_NAME(x) MAKE_FN_NAME(x)
void FN_NAME(NUM) ()
{
  int x, d;
  #pragma omp target map(from: x, d)
    {
      x = NUM;
      d = omp_is_initial_device ();
    }
  printf ("%s:\t%s ()\tx = %d\n", d ? "HOST" : "TARGET", __FUNCTION__, x);
  if (x != NUM)
    printf ("--------^\n");
}

$ gcc -DNUM=1 -c -flto test.c -o obj1.o
$ gcc -DNUM=2 -c -fopenmp test.c -o obj2.o
$ gcc -DNUM=3 -c test.c -o obj3.o
$ gcc -DNUM=4 -c -flto -fopenmp test.c -o obj4.o
$ gcc -c main.c -o main.o
$ gcc -fopenmp obj1.o obj2.o obj3.o obj4.o main.o && ./a.out
$ gcc -fopenmp obj2.o obj3.o obj4.o obj1.o main.o && ./a.out
$ gcc -fopenmp obj3.o obj1.o obj2.o obj4.o main.o && ./a.out

gcc/
	PR driver/68463
	* config/i386/intelmic-mkoffload.c (generate_target_descr_file): Don't
	define __offload_func_table and __offload_var_table.
	(generate_target_offloadend_file): Remove function.
	(prepare_target_image): Don't call generate_target_offloadend_file.
	* lto-wrapper.c (offloadbegin, offloadend): Remove static vars.
	(offload_objects_file_name): New static var.
	(tool_cleanup): Remove offload_objects_file_name file.
	(find_offloadbeginend): Rename to ...
	(find_crtoffload): ... this.  Locate crtoffload.o instead of
	crtoffloadbegin.o and crtoffloadend.o.
	(run_gcc): Remove offload_argc and offload_argv.
	Get offload_objects_file_name from -foffload-objects=... option.
	Read names of object files with offload from this file, pass them to
	compile_images_for_offload_targets.  Call find_crtoffload instead of
	find_offloadbeginend.  Don't give offload files to the linker when LTO
	is disabled, because now they're not claimed, therefore not discarded.
libgcc/
	PR driver/68463
	* Makefile.in (crtoffloadbegin$(objext)): Remove rule.
	(crtoffloadend$(objext)): Likewise.
	(crtoffload$(objext), link-offload-tables.x): New rules.
	* configure: Regenerate.
	* configure.ac (extra_parts): Add link-offload-tables.x if offloading is
	enabled, or if this is an accel compiler for intelmic.
	* link-offload-tables.x: New file.
	* offloadstuff.c: Do not define __offload_func_table,
	__offload_var_table, __offload_funcs_end, __offload_vars_end.
libgomp/
	PR driver/68463
	* Makefile.in: Regenerate.
	* configure: Regenerate.
	* configure.ac (link_offload_tables): New output variable.  Set to
	"%Tlink-offload-tables.x" if offloading is enabled, or if this is an
	accel compiler for intelmic.
	* libgomp.spec.in (*link_gomp): Add @link_offload_tables@.
	* testsuite/Makefile.in: Regenerate.
lto-plugin/
	PR driver/68463
	* lto-plugin.c (offload_files): Replace with ...
	(offload_files_1, offload_files_2, offload_files_3): ... this.
	(num_offload_files): Replace with ...
	(num_offload_files_1, num_offload_files_2, num_offload_files_3): ..this.
	(free_2): Adjust accordingly.
	(all_symbols_read_handler): Don't add offload files to lto_arg_ptr.
	Don't call free_1 for offload_files.  Write names of object files with
	offloading to the temporary file.  Add new option to lto_arg_ptr.
	(claim_file_handler): Don't claim file if it contains offload sections
	without LTO sections, add it to offload_files_1 or to offload_files_3.
	Add files with offload and LTO sections to offload_files_2.

Thanks,
  -- Ilya

[RFC,Offloading] Fix PR68463

Commit Message

Comments

Patch