[gomp4] Add tables generation

Ilya Verbin iverbin@gmail.com
Tue Sep 2 17:50:00 GMT 2014


Hi Bernd,

This patch allows to compile binaries with offloading without passing -flto option, and
w/o performing link-time optimizations of the host code.

How it works:
1.  If there is at least one function or global variable to offload, gcc sets flag_generate_lto.
This enables writing the bytecode produced by ipa_write_summaries into
.gnu.target_lto_* sections (.gnu.lto_* sections are not created).
Also this flag emits LTO marker (__gnu_lto_v1).
2.  This step is not changed: collect2 scans object files for the LTO marker and fills the list
of LTO objects.  If the list is not empty, it runs lto-wrapper to perform link-time recompilation.
3.  lto-wrapper compiles images for targets.  And if -flto option is absent
(lto_mode == LTO_MODE_NONE), then it just returns the list of input objects without recompilation.

One known issue -- the final binary contains temporary .gnu.target_lto_* sections.
This can be solved by adding the following linker script to the list of input files:
SECTIONS { /DISCARD/ : { *(.gnu.target_lto_*) } }
But I'm sure what is the best way to this automatically.

Bootstrap and make check passed, tests with '#pragma omp target' without -flto passed.
What do you think?

Thanks,
  -- Ilya


---
 gcc/cgraphunit.c  | 39 +++++++++++++++++++++++--------
 gcc/lto-wrapper.c | 68 +++++++++++++++++++++++++++++--------------------------
 gcc/omp-low.c     |  6 +++++
 gcc/passes.c      |  2 +-
 4 files changed, 73 insertions(+), 42 deletions(-)

diff --git a/gcc/cgraphunit.c b/gcc/cgraphunit.c
index f0c9f5c..32b35f3 100644
--- a/gcc/cgraphunit.c
+++ b/gcc/cgraphunit.c
@@ -2040,13 +2040,26 @@ output_in_order (void)
   free (nodes);
 }
 
-/* Collect all global variables with "omp declare target" attribute into
-   OFFLOAD_VARS.  It will be streamed out in ipa_write_summaries.  */
+/* Check whether there is at least one function or global variable to offload.
+   Also collect all such global variables into OFFLOAD_VARS, the functions were
+   already collected in omp-low.c.  They will be streamed out in
+   ipa_write_summaries.  */
 
-static void
-init_offload_var_table (void)
+static bool
+initialize_offload (void)
 {
+  bool have_offload = false;
+  struct cgraph_node *node;
   struct varpool_node *vnode;
+
+  FOR_EACH_DEFINED_FUNCTION (node)
+    if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (node->decl)))
+      {
+	have_offload = true;
+	break;
+      }
+
   FOR_EACH_DEFINED_VARIABLE (vnode)
     {
       if (!lookup_attribute ("omp declare target",
@@ -2054,13 +2067,17 @@ init_offload_var_table (void)
 	  || TREE_CODE (vnode->decl) != VAR_DECL
 	  || DECL_SIZE (vnode->decl) == 0)
 	continue;
+      have_offload = true;
       vec_safe_push (offload_vars, vnode->decl);
     }
+
+  return have_offload;
 }
 
 static void
 ipa_passes (void)
 {
+  bool have_offload = false;
   gcc::pass_manager *passes = g->get_passes ();
 
   set_cfun (NULL);
@@ -2068,6 +2085,14 @@ ipa_passes (void)
   gimple_register_cfg_hooks ();
   bitmap_obstack_initialize (NULL);
 
+  if (!in_lto_p && (flag_openacc || flag_openmp))
+    {
+      have_offload = initialize_offload ();
+      /* OpenACC / OpenMP offloading requires LTO infrastructure.  */
+      if (have_offload)
+	flag_generate_lto = 1;
+    }
+
   invoke_plugin_callbacks (PLUGIN_ALL_IPA_PASSES_START, NULL);
 
   if (!in_lto_p)
@@ -2108,11 +2133,7 @@ ipa_passes (void)
 
   if (!in_lto_p)
     {
-      init_offload_var_table ();
-
-      if ((flag_openacc || flag_openmp)
-	  && !(vec_safe_is_empty (offload_funcs)
-	       && vec_safe_is_empty (offload_vars)))
+      if (have_offload)
 	{
 	  section_name_prefix = OMP_SECTION_NAME_PREFIX;
 	  ipa_write_summaries (true);
diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c
index 80d10f3..e9245f1 100644
--- a/gcc/lto-wrapper.c
+++ b/gcc/lto-wrapper.c
@@ -668,6 +668,11 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
+      /* We may choose not to write out this .opts section in the future.  In
+	 that case we'll have to use something else to look for.  */
+      if (simple_object_find_section (sobj, OMP_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))
 	{
@@ -675,11 +680,6 @@ run_gcc (unsigned argc, char *argv[])
 	  close (fd);
 	  continue;
 	}
-      /* We may choose not to write out this .opts section in the future.  In
-	 that case we'll have to use something else to look for.  */
-      if (simple_object_find_section (sobj, OMP_SECTION_NAME_PREFIX "." "opts",
-				      &offset, &length, &errmsg, &err))
-	have_offload = true;
       lseek (fd, file_offset + offset, SEEK_SET);
       data = (char *)xmalloc (length);
       read (fd, data, length);
@@ -871,7 +871,31 @@ 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 (lto_mode == LTO_MODE_LTO)
+  if (have_offload)
+    {
+      compile_images_for_openmp_targets (argc, argv);
+      if (offload_names)
+	{
+	  find_ompbeginend ();
+	  for (i = 0; offload_names[i]; i++)
+	    printf ("%s\n", offload_names[i]);
+	  free_array_of_ptrs ((void **) offload_names, i);
+	}
+    }
+
+  if (ompbegin)
+    printf ("%s\n", ompbegin);
+
+  if (lto_mode == LTO_MODE_NONE)
+    {
+      /* If we are in lto-wrapper, but -flto option is absent, it means that
+	 there is no need to perform a link-time recompilation, i.e. lto-wrapper
+	 is used only for compiling offload images.  */
+      for (i = 1; i < argc; ++i)
+	printf ("%s\n", argv[i]);
+      goto finish;
+    }
+  else if (lto_mode == LTO_MODE_LTO)
     {
       flto_out = make_temp_file (".lto.o");
       if (linker_output)
@@ -879,7 +903,7 @@ run_gcc (unsigned argc, char *argv[])
       obstack_ptr_grow (&argv_obstack, "-o");
       obstack_ptr_grow (&argv_obstack, flto_out);
     }
-  else 
+  else if (lto_mode == LTO_MODE_WHOPR)
     {
       const char *list_option = "-fltrans-output-list=";
       size_t list_option_len = strlen (list_option);
@@ -939,7 +963,7 @@ run_gcc (unsigned argc, char *argv[])
       free (flto_out);
       flto_out = NULL;
     }
-  else
+  else if (lto_mode == LTO_MODE_WHOPR)
     {
       FILE *stream = fopen (ltrans_output_file, "r");
       FILE *mstream = NULL;
@@ -1084,25 +1108,6 @@ cont:
 	  for (i = 0; i < nr; ++i)
 	    maybe_unlink (input_names[i]);
 	}
-      if (have_offload)
-	{
-	  compile_images_for_openmp_targets (argc, argv);
-	  if (offload_names)
-	    {
-	      find_ompbeginend ();
-	      for (i = 0; offload_names[i]; i++)
-		{
-		  fputs (offload_names[i], stdout);
-		  putc ('\n', stdout);
-		}
-	      free_array_of_ptrs ((void **)offload_names, i);
-	    }
-	}
-      if (ompbegin)
-	{
-	  fputs (ompbegin, stdout);
-	  putc ('\n', stdout);
-	}
 
       for (i = 0; i < nr; ++i)
 	{
@@ -1110,11 +1115,6 @@ cont:
 	  putc ('\n', stdout);
 	  free (input_names[i]);
 	}
-      if (ompend)
-	{
-	  fputs (ompend, stdout);
-	  putc ('\n', stdout);
-	}
       nr = 0;
       free (output_names);
       free (input_names);
@@ -1122,6 +1122,10 @@ cont:
       obstack_free (&env_obstack, NULL);
     }
 
+finish:
+  if (ompend)
+    printf ("%s\n", ompend);
+
   obstack_free (&argv_obstack, NULL);
 }
 
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 1ad98ab..9289031 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -13771,6 +13771,12 @@ omp_finish_file (void)
 				   get_identifier (".omp_var_table"),
 				   vars_decl_type);
       TREE_STATIC (funcs_decl) = TREE_STATIC (vars_decl) = 1;
+      /* Do not align tables more than TYPE_ALIGN (pointer_sized_int_node),
+	 otherwise a joint table in a binary will contain padding between
+	 tables from multiple object files.  */
+      DECL_USER_ALIGN (funcs_decl) = DECL_USER_ALIGN (vars_decl) = 1;
+      DECL_ALIGN (funcs_decl) = TYPE_ALIGN (funcs_decl_type);
+      DECL_ALIGN (vars_decl) = TYPE_ALIGN (vars_decl_type);
       DECL_INITIAL (funcs_decl) = ctor_f;
       DECL_INITIAL (vars_decl) = ctor_v;
       set_decl_section_name (funcs_decl, funcs_section_name);
diff --git a/gcc/passes.c b/gcc/passes.c
index 8172185..e776059 100644
--- a/gcc/passes.c
+++ b/gcc/passes.c
@@ -2303,7 +2303,7 @@ ipa_write_summaries (bool is_omp)
   struct cgraph_node *node;
   struct cgraph_node **order;
 
-  if (!(flag_generate_lto || flag_openacc || flag_openmp) || seen_error () )
+  if (!flag_generate_lto || seen_error ())
     return;
 
   select_what_to_dump (is_omp);
-- 
1.8.3.1



More information about the Gcc-patches mailing list