public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH 2/n] OpenMP 4.0 offloading infrastructure: LTO streaming
@ 2014-09-27 18:17 Ilya Verbin
  2014-09-29  1:10 ` Jan Hubicka
  0 siblings, 1 reply; 62+ messages in thread
From: Ilya Verbin @ 2014-09-27 18:17 UTC (permalink / raw)
  To: Jakub Jelinek, Richard Biener, Jan Hubicka, gcc-patches
  Cc: Kirill Yukhin, Ilya Tocar, Andrey Turetskiy, Bernd Schmidt,
	Thomas Schwinge

Hello,

This patch enables the streaming of LTO bytecode, needed by offload target,
using existing LTO infrastructure.  It creates new prefix for the section names
(.gnu.target_lto_) and streams out the functions and variables with "omp declare
target" attribute, including the functions for outlined '#pragma omp target'
regions.  The offload compiler (under ifdef ACCEL_COMPILER) reads and compiles
these new sections.

But I have doubts regarding the offload_lto_mode switch.  Why I added it:
The outlined target regions (say omp_fn0) contains references from the parent
functions.  And that's correct for the case when we stream out the host-side
version of omp_fn0.  But for the target version there are no parent functions,
node->used_from_other_partition gets incorrect value (always 1), and offload
compiler crashes on streaming in.

Another solution is to remain referenced_from_other_partition_p and
reachable_from_other_partition_p unchanged, then used_from_other_partition will
have incorrect value for target regions, but the offload compiler will just
ignore it.  Which approach is better?
Anyway, now it's bootstrapped and regtested on i686-linux and x86_64-linux.


2014-09-27  Ilya Verbin  <ilya.verbin@intel.com>
	    Ilya Tocar  <ilya.tocar@intel.com>
	    Andrey Turetskiy  <andrey.turetskiy@intel.com>
	    Bernd Schmidt  <bernds@codesourcery.com>
gcc/
	* cgraph.h (symtab_node): Add need_dump flag.
	* cgraphunit.c: Include lto-section-names.h.
	(initialize_offload): New function.
	(ipa_passes): Initialize offload and call ipa_write_summaries if there
	is something to write to OMP_SECTION_NAME_PREFIX sections.
	(symbol_table::compile): Call lto_streamer_hooks_init under flag_openmp.
	* ipa-inline-analysis.c (inline_generate_summary): Do not exit under
	flag_openmp.
	(inline_free_summary): Always remove hooks.
	* lto-cgraph.c (lto_set_symtab_encoder_in_partition): Exit if there is
	no need to encode the node.
	(referenced_from_other_partition_p, reachable_from_other_partition_p):
	Ignore references from non-target functions to target functions if we
	are streaming out target-side bytecode (offload lto mode).
	(select_what_to_dump): New function.
	* lto-section-names.h (OMP_SECTION_NAME_PREFIX): Define.
	(section_name_prefix): Declare.
	* lto-streamer.c (offload_lto_mode): New variable.
	(section_name_prefix): New variable.
	(lto_get_section_name): Use section_name_prefix instead of
	LTO_SECTION_NAME_PREFIX.
	* lto-streamer.h (select_what_to_dump): Declare.
	(offload_lto_mode): Declare.
	* omp-low.c (is_targetreg_ctx): New function.
	(create_omp_child_function, check_omp_nesting_restrictions): Use it.
	(expand_omp_target): Set mark_force_output for the target functions.
	(lower_omp_critical): Add target attribute for omp critical symbol.
	* passes.c (ipa_write_summaries): Call select_what_to_dump.
gcc/lto/
	* lto-object.c (lto_obj_add_section): Use section_name_prefix instead of
	LTO_SECTION_NAME_PREFIX.
	* lto-partition.c (add_symbol_to_partition_1): Always set
	node->need_dump to true.
	(lto_promote_cross_file_statics): Call select_what_to_dump.
	* lto.c (lto_section_with_id): Use section_name_prefix instead of
	LTO_SECTION_NAME_PREFIX.
	(read_cgraph_and_symbols): Read OMP_SECTION_NAME_PREFIX sections, if
	being built as an offload compiler.

Thanks,
  -- Ilya

---

diff --git a/gcc/cgraph.h b/gcc/cgraph.h
index 7481906..9ab970d 100644
--- a/gcc/cgraph.h
+++ b/gcc/cgraph.h
@@ -444,6 +444,11 @@ public:
   /* Set when init priority is set.  */
   unsigned in_init_priority_hash : 1;
 
+  /* Set when symbol needs to be dumped into LTO bytecode for LTO,
+     or in pragma omp target case, for separate compilation targeting
+     a different architecture.  */
+  unsigned need_dump : 1;
+
 
   /* Ordering of all symtab entries.  */
   int order;
diff --git a/gcc/cgraphunit.c b/gcc/cgraphunit.c
index b854e4b..4ab4c57 100644
--- a/gcc/cgraphunit.c
+++ b/gcc/cgraphunit.c
@@ -211,6 +211,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "tree-nested.h"
 #include "gimplify.h"
 #include "dbgcnt.h"
+#include "lto-section-names.h"
 
 /* Queue of cgraph nodes scheduled to be added into cgraph.  This is a
    secondary queue used during optimization to accommodate passes that
@@ -1994,9 +1995,40 @@ output_in_order (bool no_reorder)
   free (nodes);
 }
 
+/* Check whether there is at least one function or global variable to offload.
+   */
+
+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",
+			     DECL_ATTRIBUTES (vnode->decl))
+	  || TREE_CODE (vnode->decl) != VAR_DECL
+	  || DECL_SIZE (vnode->decl) == 0)
+	continue;
+      have_offload = true;
+    }
+
+  return have_offload;
+}
+
 static void
 ipa_passes (void)
 {
+  bool have_offload = false;
   gcc::pass_manager *passes = g->get_passes ();
 
   set_cfun (NULL);
@@ -2004,6 +2036,14 @@ ipa_passes (void)
   gimple_register_cfg_hooks ();
   bitmap_obstack_initialize (NULL);
 
+  if (!in_lto_p && flag_openmp)
+    {
+      have_offload = initialize_offload ();
+      /* 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)
@@ -2041,7 +2081,20 @@ ipa_passes (void)
     targetm.asm_out.lto_start ();
 
   if (!in_lto_p)
-    ipa_write_summaries ();
+    {
+      if (have_offload)
+	{
+	  offload_lto_mode = true;
+	  section_name_prefix = OMP_SECTION_NAME_PREFIX;
+	  ipa_write_summaries ();
+	}
+      if (flag_lto)
+	{
+	  offload_lto_mode = false;
+	  section_name_prefix = LTO_SECTION_NAME_PREFIX;
+	  ipa_write_summaries ();
+	}
+    }
 
   if (flag_generate_lto)
     targetm.asm_out.lto_end ();
@@ -2122,7 +2175,7 @@ symbol_table::compile (void)
   state = IPA;
 
   /* If LTO is enabled, initialize the streamer hooks needed by GIMPLE.  */
-  if (flag_lto)
+  if (flag_lto || flag_openmp)
     lto_streamer_hooks_init ();
 
   /* Don't run the IPA passes if there was any error or sorry messages.  */
diff --git a/gcc/ipa-inline-analysis.c b/gcc/ipa-inline-analysis.c
index 38f56d2..076a1e8 100644
--- a/gcc/ipa-inline-analysis.c
+++ b/gcc/ipa-inline-analysis.c
@@ -4010,7 +4010,7 @@ inline_generate_summary (void)
 
   /* When not optimizing, do not bother to analyze.  Inlining is still done
      because edge redirection needs to happen there.  */
-  if (!optimize && !flag_lto && !flag_wpa)
+  if (!optimize && !flag_lto && !flag_wpa && !flag_openmp)
     return;
 
   function_insertion_hook_holder =
@@ -4325,11 +4325,6 @@ void
 inline_free_summary (void)
 {
   struct cgraph_node *node;
-  if (!inline_edge_summary_vec.exists ())
-    return;
-  FOR_EACH_DEFINED_FUNCTION (node)
-    if (!node->alias)
-      reset_inline_summary (node);
   if (function_insertion_hook_holder)
     symtab->remove_cgraph_insertion_hook (function_insertion_hook_holder);
   function_insertion_hook_holder = NULL;
@@ -4345,6 +4340,11 @@ inline_free_summary (void)
   if (edge_duplication_hook_holder)
     symtab->remove_edge_duplication_hook (edge_duplication_hook_holder);
   edge_duplication_hook_holder = NULL;
+  if (!inline_edge_summary_vec.exists ())
+    return;
+  FOR_EACH_DEFINED_FUNCTION (node)
+    if (!node->alias)
+      reset_inline_summary (node);
   vec_free (inline_summary_vec);
   inline_edge_summary_vec.release ();
   if (edge_predicate_pool)
diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
index 0584946..78b7fc8 100644
--- a/gcc/lto-cgraph.c
+++ b/gcc/lto-cgraph.c
@@ -239,6 +239,9 @@ void
 lto_set_symtab_encoder_in_partition (lto_symtab_encoder_t encoder,
 				     symtab_node *node)
 {
+  /* Ignore not needed nodes.  */
+  if (!node->need_dump)
+    return;
   int index = lto_symtab_encoder_encode (encoder, node);
   encoder->nodes[index].in_partition = true;
 }
@@ -321,6 +324,12 @@ referenced_from_other_partition_p (symtab_node *node, lto_symtab_encoder_t encod
 
   for (i = 0; node->iterate_referring (i, ref); i++)
     {
+      /* Ignore references from non-target functions in offload lto mode.  */
+      if (offload_lto_mode
+	  && !lookup_attribute ("omp declare target",
+				DECL_ATTRIBUTES (ref->referring->decl)))
+	continue;
+
       if (ref->referring->in_other_partition
           || !lto_symtab_encoder_in_partition_p (encoder, ref->referring))
 	return true;
@@ -339,9 +348,17 @@ reachable_from_other_partition_p (struct cgraph_node *node, lto_symtab_encoder_t
   if (node->global.inlined_to)
     return false;
   for (e = node->callers; e; e = e->next_caller)
-    if (e->caller->in_other_partition
-	|| !lto_symtab_encoder_in_partition_p (encoder, e->caller))
-      return true;
+    {
+      /* Ignore references from non-target functions in offload lto mode.  */
+      if (offload_lto_mode
+	  && !lookup_attribute ("omp declare target",
+				DECL_ATTRIBUTES (e->caller->decl)))
+	continue;
+
+      if (e->caller->in_other_partition
+	  || !lto_symtab_encoder_in_partition_p (encoder, e->caller))
+	return true;
+    }
   return false;
 }
 
@@ -802,6 +819,18 @@ create_references (lto_symtab_encoder_t encoder, symtab_node *node)
       lto_symtab_encoder_encode (encoder, ref->referred);
 }
 
+/* Select what needs to be streamed out.  In regular lto mode stream everything.
+   In offload lto mode stream only stuff marked with an attribute.  */
+void
+select_what_to_dump (void)
+{
+  struct symtab_node *snode;
+  FOR_EACH_SYMBOL (snode)
+    snode->need_dump = !offload_lto_mode
+		       || lookup_attribute ("omp declare target",
+					    DECL_ATTRIBUTES (snode->decl));
+}
+
 /* Find all symbols we want to stream into given partition and insert them
    to encoders.
 
diff --git a/gcc/lto-section-names.h b/gcc/lto-section-names.h
index cb75230..06d2caf 100644
--- a/gcc/lto-section-names.h
+++ b/gcc/lto-section-names.h
@@ -25,6 +25,11 @@ along with GCC; see the file COPYING3.  If not see
    name for the functions and static_initializers.  For other types of
    sections a '.' and the section type are appended.  */
 #define LTO_SECTION_NAME_PREFIX ".gnu.lto_"
+#define OMP_SECTION_NAME_PREFIX ".gnu.target_lto_"
+
+/* Can be either OMP_SECTION_NAME_PREFIX when we stream 'pragma omp target'
+   stuff, or LTO_SECTION_NAME_PREFIX for LTO case.  */
+extern const char *section_name_prefix;
 
 /* Segment name for LTO sections.  This is only used for Mach-O.  */
 
diff --git a/gcc/lto-streamer.c b/gcc/lto-streamer.c
index 3480723..95232f9 100644
--- a/gcc/lto-streamer.c
+++ b/gcc/lto-streamer.c
@@ -48,6 +48,8 @@ struct lto_stats_d lto_stats;
 static bitmap_obstack lto_obstack;
 static bool lto_obstack_initialized;
 
+bool offload_lto_mode = false;
+const char *section_name_prefix = LTO_SECTION_NAME_PREFIX;
 
 /* Return a string representing LTO tag TAG.  */
 
@@ -177,7 +179,7 @@ lto_get_section_name (int section_type, const char *name, struct lto_file_decl_d
     sprintf (post, "." HOST_WIDE_INT_PRINT_HEX_PURE, f->id);
   else
     sprintf (post, "." HOST_WIDE_INT_PRINT_HEX_PURE, get_random_seed (false)); 
-  return concat (LTO_SECTION_NAME_PREFIX, sep, add, post, NULL);
+  return concat (section_name_prefix, sep, add, post, NULL);
 }
 
 
diff --git a/gcc/lto-streamer.h b/gcc/lto-streamer.h
index 4bec969..0016eef 100644
--- a/gcc/lto-streamer.h
+++ b/gcc/lto-streamer.h
@@ -831,6 +831,7 @@ bool referenced_from_this_partition_p (symtab_node *,
 bool reachable_from_this_partition_p (struct cgraph_node *,
 				      lto_symtab_encoder_t);
 lto_symtab_encoder_t compute_ltrans_boundary (lto_symtab_encoder_t encoder);
+void select_what_to_dump (void);
 
 
 /* In lto-symtab.c.  */
@@ -846,6 +847,9 @@ extern void lto_write_options (void);
 /* Statistics gathered during LTO, WPA and LTRANS.  */
 extern struct lto_stats_d lto_stats;
 
+/* Regular or offload mode of LTO.  */
+extern bool offload_lto_mode;
+
 /* Section names corresponding to the values of enum lto_section_type.  */
 extern const char *lto_section_name[];
 
diff --git a/gcc/lto/lto-object.c b/gcc/lto/lto-object.c
index 323f7b2..4ee752f 100644
--- a/gcc/lto/lto-object.c
+++ b/gcc/lto/lto-object.c
@@ -230,8 +230,7 @@ lto_obj_add_section (void *data, const char *name, off_t offset,
   void **slot;
   struct lto_section_list *list = loasd->list;
 
-  if (strncmp (name, LTO_SECTION_NAME_PREFIX,
-	       strlen (LTO_SECTION_NAME_PREFIX)) != 0)
+  if (strncmp (name, section_name_prefix, strlen (section_name_prefix)))
     return 1;
 
   new_name = xstrdup (name);
diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
index 0451a66..332562f 100644
--- a/gcc/lto/lto-partition.c
+++ b/gcc/lto/lto-partition.c
@@ -134,6 +134,7 @@ add_symbol_to_partition_1 (ltrans_partition part, symtab_node *node)
   gcc_assert (c != SYMBOL_EXTERNAL
 	      && (c == SYMBOL_DUPLICATE || !symbol_partitioned_p (node)));
 
+  node->need_dump = true;
   lto_set_symtab_encoder_in_partition (part->encoder, node);
 
   if (symbol_partitioned_p (node))
@@ -920,6 +921,8 @@ lto_promote_cross_file_statics (void)
 
   gcc_assert (flag_wpa);
 
+  select_what_to_dump ();
+
   /* First compute boundaries.  */
   n_sets = ltrans_partitions.length ();
   for (i = 0; i < n_sets; i++)
diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index 6cbb178..f23d997 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -2125,7 +2125,7 @@ lto_section_with_id (const char *name, unsigned HOST_WIDE_INT *id)
 {
   const char *s;
 
-  if (strncmp (name, LTO_SECTION_NAME_PREFIX, strlen (LTO_SECTION_NAME_PREFIX)))
+  if (strncmp (name, section_name_prefix, strlen (section_name_prefix)))
     return 0;
   s = strrchr (name, '.');
   return s && sscanf (s, "." HOST_WIDE_INT_PRINT_HEX_PURE, id) == 1;
@@ -2899,6 +2899,10 @@ read_cgraph_and_symbols (unsigned nfiles, const char **fnames)
 
   timevar_push (TV_IPA_LTO_DECL_IN);
 
+#ifdef ACCEL_COMPILER
+    section_name_prefix = OMP_SECTION_NAME_PREFIX;
+#endif
+
   real_file_decl_data
     = decl_data = ggc_cleared_vec_alloc<lto_file_decl_data_ptr> (nfiles + 1);
   real_file_count = nfiles;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 82651ea..7d587b3 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -257,6 +257,16 @@ is_parallel_ctx (omp_context *ctx)
 }
 
 
+/* Return true if CTX is for an omp target region.  */
+
+static inline bool
+is_targetreg_ctx (omp_context *ctx)
+{
+  return gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+	 && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION;
+}
+
+
 /* Return true if CTX is for an omp task.  */
 
 static inline bool
@@ -1930,9 +1940,7 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
     {
       omp_context *octx;
       for (octx = ctx; octx; octx = octx->outer)
-	if (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET
-	    && gimple_omp_target_kind (octx->stmt)
-	       == GF_OMP_TARGET_KIND_REGION)
+	if (is_targetreg_ctx (octx))
 	  {
 	    target_p = true;
 	    break;
@@ -2588,8 +2596,7 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
       break;
     case GIMPLE_OMP_TARGET:
       for (; ctx != NULL; ctx = ctx->outer)
-	if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
-	    && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION)
+	if (is_targetreg_ctx (ctx))
 	  {
 	    const char *name;
 	    switch (gimple_omp_target_kind (stmt))
@@ -8206,6 +8213,7 @@ expand_omp_target (struct omp_region *region)
   if (kind == GF_OMP_TARGET_KIND_REGION)
     {
       unsigned srcidx, dstidx, num;
+      struct cgraph_node *node;
 
       /* If the target region needs data sent from the parent
 	 function, then the very first statement (except possible
@@ -8337,6 +8345,11 @@ expand_omp_target (struct omp_region *region)
       push_cfun (child_cfun);
       cgraph_edge::rebuild_edges ();
 
+      /* Prevent IPA from removing child_fn as unreachable, since there are no
+	 refs from the parent function to the target side child_fn.  */
+      node = cgraph_node::get (child_fn);
+      node->mark_force_output ();
+
       /* Some EH regions might become dead, see PR34608.  If
 	 pass_cleanup_cfg isn't the first pass to happen with the
 	 new child, these dead EH edges might cause problems.
@@ -9207,6 +9220,19 @@ lower_omp_critical (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  DECL_COMMON (decl) = 1;
 	  DECL_ARTIFICIAL (decl) = 1;
 	  DECL_IGNORED_P (decl) = 1;
+
+	  /* If '#pragma omp critical' is inside target region, the symbol must
+	     have an 'omp declare target' attribute.  */
+	  omp_context *octx;
+	  for (octx = ctx->outer; octx; octx = octx->outer)
+	    if (is_targetreg_ctx (octx))
+	      {
+		DECL_ATTRIBUTES (decl)
+		  = tree_cons (get_identifier ("omp declare target"),
+			       NULL_TREE, DECL_ATTRIBUTES (decl));
+		break;
+	      }
+
 	  varpool_node::finalize_decl (decl);
 
 	  splay_tree_insert (critical_name_mutexes, (splay_tree_key) name,
diff --git a/gcc/passes.c b/gcc/passes.c
index 5001c3d..d63c913 100644
--- a/gcc/passes.c
+++ b/gcc/passes.c
@@ -2308,6 +2308,8 @@ ipa_write_summaries (void)
   if (!flag_generate_lto || seen_error ())
     return;
 
+  select_what_to_dump ();
+
   encoder = lto_symtab_encoder_new (false);
 
   /* Create the callgraph set in the same order used in
-- 
1.7.1

^ permalink raw reply	[flat|nested] 62+ messages in thread
* Pass -foffload targets from driver to libgomp at link time
@ 2015-08-27 20:58 Joseph Myers
  2015-09-03 14:58 ` Ping " Joseph Myers
  0 siblings, 1 reply; 62+ messages in thread
From: Joseph Myers @ 2015-08-27 20:58 UTC (permalink / raw)
  To: gcc-patches; +Cc: thomas, Nathan Sidwell, jakub

This patch, a version of
<https://gcc.gnu.org/ml/gcc-patches/2015-08/msg01264.html> cleaned up
for trunk, arranges for the -foffload= targets specified at link time
to be passed to libgomp via a constructor function generated by the
driver.

In this patch, I've tried to remove all the miscellaneous cleanups in
the gomp-4_0-branch version that didn't appear to be necessarily
required as part of passing -foffload from the driver to libgomp.
Thus, care should be taken when next merging from trunk to
gomp-4_0-branch not to lose those cleanups where patch conflicts
arise, where the cleanups are still desired for merging to trunk
separately.  It's possible I missed some such changes; thus, this
patch should be reviewed carefully to make sure there isn't anything
unrelated mixed in.

This patch uses GOMP_4.0.2 as the symbol version for the new function
GOMP_set_offload_targets (where the gomp-4_0-branch patch had
GOACC_2.0.GOMP_4_BRANCH).  I hope this is the correct version for a
GOMP_* function that is new in GCC 6.

Tested with no regressions for x86_64-none-linux-gnu, offloading to
nvptx-none; 24 libgomp test FAILs start to pass with the patch.  OK to
commit?

gcc:
2015-08-27  Thomas Schwinge  <thomas@codesourcery.com>
	    Joseph Myers  <joseph@codesourcery.com>

	* gcc.c (offload_targets): Update comment.
	(add_omp_infile_spec_func, spec_lang_mask_accept): New.
	(driver_self_specs) [ENABLE_OFFLOADING]: Add spec to use
	%:add-omp-infile().
	(static_spec_functions): Add add-omp-infile.
	(struct switchstr): Add lang_mask field.
	(struct infile): Add lang_mask field.
	(add_infile, save_switch, do_spec): Add lang_mask argument.
	(driver_unknown_option_callback, driver_wrong_lang_callback)
	(driver_handle_option, process_command, do_self_spec)
	(driver::do_spec_on_infiles): All callers changed.
	(process_command): Call handle_foffload_option (OFFLOAD_TARGETS)
	if no offload target specified.
	(give_switch): Check languages of switch against
	spec_lang_mask_accept.
	(driver::maybe_putenv_OFFLOAD_TARGETS): Do not use intermediate
	targets variable.
	* gcc.h (do_spec): Update prototype.

gcc/fortran:
2015-08-27  Joseph Myers  <joseph@codesourcery.com>

	* gfortranspec.c (lang_specific_pre_link): Update call to do_spec.

gcc/java:
2015-08-27  Joseph Myers  <joseph@codesourcery.com>

	* jvspec.c (lang_specific_pre_link): Update call to do_spec.

libgomp:
2015-08-27  Thomas Schwinge  <thomas@codesourcery.com>
	    Joseph Myers  <joseph@codesourcery.com>

	* plugin/configfrag.ac (tgt_name): Do not set.
	(offload_targets): Separate with colons not commas.
	* config.h.in, configure: Regenerate.
	* libgomp.map (GOMP_4.0.2): Add GOMP_set_offload_targets.
	* libgomp_g.h (GOMP_set_offload_targets): New prototype.
	* target.c (offload_target_to_plugin_name, gomp_offload_targets)
	(gomp_offload_targets_init, GOMP_set_offload_targets): New.
	(gomp_target_init): Use gomp_offload_targets instead of
	OFFLOAD_TARGETS.  Handle and rewrite colon-separated string.
	* testsuite/lib/libgomp.exp: Expect offload targets to be
	colon-separated.  Adjust matching of offload targets.
	(libgomp_init)
	(check_effective_target_openacc_nvidia_accel_supported)
	(check_effective_target_openacc_host_selected): Adjust checks of
	offload target names.
	* testsuite/libgomp.oacc-c++/c++.exp: Adjust set of offload
	targets.  Use -foffload=.
	* testsuite/libgomp.oacc-c/c.exp: Adjust set of offload targets.
	Use -foffload=.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Adjust set of
	offload targets.  Use -foffload=.

Index: libgomp/config.h.in
===================================================================
--- libgomp/config.h.in	(revision 227194)
+++ libgomp/config.h.in	(working copy)
@@ -95,7 +95,7 @@
    */
 #undef LT_OBJDIR
 
-/* Define to hold the list of target names suitable for offloading. */
+/* Define to hold the list of offload targets, separated by colons. */
 #undef OFFLOAD_TARGETS
 
 /* Name of package */
Index: libgomp/target.c
===================================================================
--- libgomp/target.c	(revision 227194)
+++ libgomp/target.c	(working copy)
@@ -1209,6 +1209,41 @@ gomp_load_plugin_for_device (struct gomp_device_de
   return 0;
 }
 
+/* Return the corresponding plugin name for the offload target name
+   OFFLOAD_TARGET.  */
+
+static const char *
+offload_target_to_plugin_name (const char *offload_target)
+{
+  if (strstr (offload_target, "-intelmic") != NULL)
+    return "intelmic";
+  if (strncmp (offload_target, "nvptx", 5) == 0)
+    return "nvptx";
+  gomp_fatal ("Unknown offload target: %s", offload_target);
+}
+
+/* List of offload targets, separated by colon.  Defaults to the list
+   determined when configuring libgomp.  */
+static const char *gomp_offload_targets = OFFLOAD_TARGETS;
+static bool gomp_offload_targets_init = false;
+
+/* Override the list of offload targets with OFFLOAD_TARGETS, the set
+   passed to the compiler at link time.  This must be called early,
+   and only once.  */
+
+void
+GOMP_set_offload_targets (const char *offload_targets)
+{
+  gomp_debug (0, "%s (\"%s\")\n", __FUNCTION__, offload_targets);
+
+  /* Make sure this gets called early.  */
+  assert (gomp_is_initialized == PTHREAD_ONCE_INIT);
+  /* Make sure this only gets called once.  */
+  assert (!gomp_offload_targets_init);
+  gomp_offload_targets_init = true;
+  gomp_offload_targets = offload_targets;
+}
+
 /* This function initializes the runtime needed for offloading.
    It parses the list of offload targets and tries to load the plugins for
    these targets.  On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP
@@ -1228,26 +1263,45 @@ gomp_target_init (void)
   num_devices = 0;
   devices = NULL;
 
-  cur = OFFLOAD_TARGETS;
+  cur = gomp_offload_targets;
   if (*cur)
     do
       {
 	struct gomp_device_descr current_device;
 
-	next = strchr (cur, ',');
-
-	plugin_name = (char *) malloc (1 + (next ? next - cur : strlen (cur))
-				       + strlen (prefix) + strlen (suffix));
+	next = strchr (cur, ':');
+	size_t prefix_len = strlen (prefix);
+	size_t cur_len = next ? next - cur : strlen (cur);
+	size_t suffix_len = strlen (suffix);
+	plugin_name = (char *) malloc (prefix_len
+				       + cur_len
+				       + suffix_len
+				       + 1);
 	if (!plugin_name)
 	  {
 	    num_devices = 0;
 	    break;
 	  }
+	memcpy (plugin_name, prefix, prefix_len);
+	memcpy (plugin_name + prefix_len, cur, cur_len);
+	/* NUL-terminate the string here...  */
+	plugin_name[prefix_len + cur_len] = '\0';
+	/* ..., so that we can then use it to translate the offload target to
+	   the plugin name...  */
+	const char *cur_plugin_name
+	  = offload_target_to_plugin_name (plugin_name
+					   + prefix_len);
+	size_t cur_plugin_name_len = strlen (cur_plugin_name);
+	assert (cur_plugin_name_len <= cur_len);
+	/* ..., and then rewrite it.  */
+	memcpy (plugin_name + prefix_len,
+		cur_plugin_name, cur_plugin_name_len);
+	memcpy (plugin_name + prefix_len + cur_plugin_name_len,
+		suffix, suffix_len);
+	plugin_name[prefix_len
+		    + cur_plugin_name_len
+		    + suffix_len] = '\0';
 
-	strcpy (plugin_name, prefix);
-	strncat (plugin_name, cur, next ? next - cur : strlen (cur));
-	strcat (plugin_name, suffix);
-
 	if (gomp_load_plugin_for_device (&current_device, plugin_name))
 	  {
 	    new_num_devices = current_device.get_num_devices_func ();
Index: libgomp/configure
===================================================================
--- libgomp/configure	(revision 227194)
+++ libgomp/configure	(working copy)
@@ -15236,10 +15236,8 @@ if test x"$enable_offload_targets" != x; then
     tgt=`echo $tgt | sed 's/=.*//'`
     case $tgt in
       *-intelmic-* | *-intelmicemul-*)
-	tgt_name=intelmic
 	;;
       nvptx*)
-        tgt_name=nvptx
 	PLUGIN_NVPTX=$tgt
 	PLUGIN_NVPTX_CPPFLAGS=$CUDA_DRIVER_CPPFLAGS
 	PLUGIN_NVPTX_LDFLAGS=$CUDA_DRIVER_LDFLAGS
@@ -15282,9 +15280,9 @@ rm -f core conftest.err conftest.$ac_objext \
 	;;
     esac
     if test x"$offload_targets" = x; then
-      offload_targets=$tgt_name
+      offload_targets=$tgt
     else
-      offload_targets=$offload_targets,$tgt_name
+      offload_targets=$offload_targets:$tgt
     fi
     if test x"$tgt_dir" != x; then
       offload_additional_options="$offload_additional_options -B$tgt_dir/libexec/gcc/\$(target_alias)/\$(gcc_version) -B$tgt_dir/bin"
Index: libgomp/libgomp_g.h
===================================================================
--- libgomp/libgomp_g.h	(revision 227194)
+++ libgomp/libgomp_g.h	(working copy)
@@ -206,6 +206,7 @@ extern void GOMP_single_copy_end (void *);
 
 /* target.c */
 
+extern void GOMP_set_offload_targets (const char *);
 extern void GOMP_target (int, void (*) (void *), const void *,
 			 size_t, void **, size_t *, unsigned char *);
 extern void GOMP_target_data (int, const void *,
Index: libgomp/plugin/configfrag.ac
===================================================================
--- libgomp/plugin/configfrag.ac	(revision 227194)
+++ libgomp/plugin/configfrag.ac	(working copy)
@@ -92,10 +92,8 @@ if test x"$enable_offload_targets" != x; then
     tgt=`echo $tgt | sed 's/=.*//'`
     case $tgt in
       *-intelmic-* | *-intelmicemul-*)
-	tgt_name=intelmic
 	;;
       nvptx*)
-        tgt_name=nvptx
 	PLUGIN_NVPTX=$tgt
 	PLUGIN_NVPTX_CPPFLAGS=$CUDA_DRIVER_CPPFLAGS
 	PLUGIN_NVPTX_LDFLAGS=$CUDA_DRIVER_LDFLAGS
@@ -127,9 +125,9 @@ if test x"$enable_offload_targets" != x; then
 	;;
     esac
     if test x"$offload_targets" = x; then
-      offload_targets=$tgt_name
+      offload_targets=$tgt
     else
-      offload_targets=$offload_targets,$tgt_name
+      offload_targets=$offload_targets:$tgt
     fi
     if test x"$tgt_dir" != x; then
       offload_additional_options="$offload_additional_options -B$tgt_dir/libexec/gcc/\$(target_alias)/\$(gcc_version) -B$tgt_dir/bin"
@@ -141,7 +139,7 @@ if test x"$enable_offload_targets" != x; then
   done
 fi
 AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets",
-  [Define to hold the list of target names suitable for offloading.])
+  [Define to hold the list of offload targets, separated by colons.])
 AM_CONDITIONAL([PLUGIN_NVPTX], [test $PLUGIN_NVPTX = 1])
 AC_DEFINE_UNQUOTED([PLUGIN_NVPTX], [$PLUGIN_NVPTX],
   [Define to 1 if the NVIDIA plugin is built, 0 if not.])
Index: libgomp/libgomp.map
===================================================================
--- libgomp/libgomp.map	(revision 227194)
+++ libgomp/libgomp.map	(working copy)
@@ -238,6 +238,7 @@ GOMP_4.0.2 {
   global:
 	GOMP_offload_register_ver;
 	GOMP_offload_unregister_ver;
+	GOMP_set_offload_targets;
 } GOMP_4.0.1;
 
 OACC_2.0 {
Index: libgomp/testsuite/libgomp.oacc-c++/c++.exp
===================================================================
--- libgomp/testsuite/libgomp.oacc-c++/c++.exp	(revision 227194)
+++ libgomp/testsuite/libgomp.oacc-c++/c++.exp	(working copy)
@@ -75,13 +75,12 @@ if { $lang_test_file_found } {
 
     # Test OpenACC with available accelerators.
     foreach offload_target_openacc $offload_targets_s_openacc {
-	set tagopt "-DACC_DEVICE_TYPE_$offload_target_openacc=1"
-
-	switch $offload_target_openacc {
-	    host {
+	switch -glob $offload_target_openacc {
+	    disable {
 		set acc_mem_shared 1
+		set tagopt "-DACC_DEVICE_TYPE_host=1"
 	    }
-	    nvidia {
+	    nvptx* {
 		if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.
 		    untested "$subdir $offload_target_openacc offloading"
@@ -95,15 +94,14 @@ if { $lang_test_file_found } {
 		lappend ALWAYS_CFLAGS "additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common"
 
 		set acc_mem_shared 0
+		set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
 	    }
 	    default {
 		set acc_mem_shared 0
 	    }
 	}
-	set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
+	set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared -foffload=$offload_target_openacc"
 
-	setenv ACC_DEVICE_TYPE $offload_target_openacc
-
 	dg-runtest $tests "$tagopt" "$libstdcxx_includes $DEFAULT_CFLAGS"
     }
 }
Index: libgomp/testsuite/lib/libgomp.exp
===================================================================
--- libgomp/testsuite/lib/libgomp.exp	(revision 227194)
+++ libgomp/testsuite/lib/libgomp.exp	(working copy)
@@ -36,24 +36,21 @@ load_gcc_lib fortran-modules.exp
 load_file libgomp-test-support.exp
 
 # Populate offload_targets_s (offloading targets separated by a space), and
-# offload_targets_s_openacc (the same, but with OpenACC names; OpenACC spells
-# some of them a little differently).
-set offload_targets_s [split $offload_targets ","]
+# offload_targets_s_openacc (those suitable for OpenACC).
+set offload_targets_s [split $offload_targets ":"]
 set offload_targets_s_openacc {}
 foreach offload_target_openacc $offload_targets_s {
-    switch $offload_target_openacc {
-	intelmic {
+    switch -glob $offload_target_openacc {
+	*-intelmic* {
 	    # Skip; will all FAIL because of missing
 	    # GOMP_OFFLOAD_CAP_OPENACC_200.
 	    continue
 	}
-	nvptx {
-	    set offload_target_openacc "nvidia"
-	}
     }
     lappend offload_targets_s_openacc "$offload_target_openacc"
 }
-lappend offload_targets_s_openacc "host"
+# Host fallback.
+lappend offload_targets_s_openacc "disable"
 
 set dg-do-what-default run
 
@@ -134,7 +131,7 @@ proc libgomp_init { args } {
     # Add liboffloadmic build directory in LD_LIBRARY_PATH to support
     # non-fallback testing for Intel MIC targets
     global offload_targets
-    if { [string match "*,intelmic,*" ",$offload_targets,"] } {
+    if { [string match "*:*-intelmic*:*" ":$offload_targets:"] } {
 	append always_ld_library_path ":${blddir}/../liboffloadmic/.libs"
 	append always_ld_library_path ":${blddir}/../liboffloadmic/plugin/.libs"
 	# libstdc++ is required by liboffloadmic
@@ -332,15 +329,14 @@ proc check_effective_target_openacc_nvidia_accel_p
 }
 
 # Return 1 if at least one nvidia board is present, and the nvidia device type
-# is selected by default by means of setting the environment variable
-# ACC_DEVICE_TYPE.
+# is selected by default.
 
 proc check_effective_target_openacc_nvidia_accel_selected { } {
     if { ![check_effective_target_openacc_nvidia_accel_present] } {
 	return 0;
     }
     global offload_target_openacc
-    if { $offload_target_openacc == "nvidia" } {
+    if { [string match "nvptx*" $offload_target_openacc] } {
         return 1;
     }
     return 0;
@@ -350,7 +346,7 @@ proc check_effective_target_openacc_nvidia_accel_s
 
 proc check_effective_target_openacc_host_selected { } {
     global offload_target_openacc
-    if { $offload_target_openacc == "host" } {
+    if { $offload_target_openacc == "disable" } {
         return 1;
     }
     return 0;
Index: libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/fortran.exp	(revision 227194)
+++ libgomp/testsuite/libgomp.oacc-fortran/fortran.exp	(working copy)
@@ -67,13 +67,12 @@ if { $lang_test_file_found } {
 
     # Test OpenACC with available accelerators.
     foreach offload_target_openacc $offload_targets_s_openacc {
-	set tagopt "-DACC_DEVICE_TYPE_$offload_target_openacc=1"
-
-	switch $offload_target_openacc {
-	    host {
+	switch -glob $offload_target_openacc {
+	    disable {
 		set acc_mem_shared 1
+		set tagopt "-DACC_DEVICE_TYPE_host=1"
 	    }
-	    nvidia {
+	    nvptx* {
 		if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.
 		    untested "$subdir $offload_target_openacc offloading"
@@ -81,15 +80,14 @@ if { $lang_test_file_found } {
 		}
 
 		set acc_mem_shared 0
+		set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
 	    }
 	    default {
 		set acc_mem_shared 0
 	    }
 	}
-	set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
+	set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared -foffload=$offload_target_openacc"
 
-	setenv ACC_DEVICE_TYPE $offload_target_openacc
-
 	# For Fortran we're doing torture testing, as Fortran has far more tests
 	# with arrays etc. that testing just -O0 or -O2 is insufficient, that is
 	# typically not the case for C/C++.
Index: libgomp/testsuite/libgomp.oacc-c/c.exp
===================================================================
--- libgomp/testsuite/libgomp.oacc-c/c.exp	(revision 227194)
+++ libgomp/testsuite/libgomp.oacc-c/c.exp	(working copy)
@@ -38,13 +38,13 @@ set_ld_library_path_env_vars
 set SAVE_ALWAYS_CFLAGS "$ALWAYS_CFLAGS"
 foreach offload_target_openacc $offload_targets_s_openacc {
     set ALWAYS_CFLAGS "$SAVE_ALWAYS_CFLAGS"
-    set tagopt "-DACC_DEVICE_TYPE_$offload_target_openacc=1"
 
-    switch $offload_target_openacc {
-	host {
+    switch -glob $offload_target_openacc {
+	disable {
 	    set acc_mem_shared 1
+	    set tagopt "-DACC_DEVICE_TYPE_host=1"
 	}
-	nvidia {
+	nvptx* {
 	    if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		# Don't bother; execution testing is going to FAIL.
 		untested "$subdir $offload_target_openacc offloading"
@@ -58,15 +58,14 @@ foreach offload_target_openacc $offload_targets_s_
 	    lappend ALWAYS_CFLAGS "additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common"
 
 	    set acc_mem_shared 0
+	    set tagopt "-DACC_DEVICE_TYPE_nvidia=1"
 	}
 	default {
 	    set acc_mem_shared 0
 	}
     }
-    set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
+    set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared -foffload=$offload_target_openacc"
 
-    setenv ACC_DEVICE_TYPE $offload_target_openacc
-
     dg-runtest $tests "$tagopt" $DEFAULT_CFLAGS
 }
 
Index: gcc/java/jvspec.c
===================================================================
--- gcc/java/jvspec.c	(revision 227194)
+++ gcc/java/jvspec.c	(working copy)
@@ -629,7 +629,7 @@ lang_specific_pre_link (void)
      class name.  Append dummy `.c' that can be stripped by set_input so %b
      is correct.  */ 
   set_input (concat (main_class_name, "main.c", NULL));
-  err = do_spec (jvgenmain_spec);
+  err = do_spec (jvgenmain_spec, 0);
   if (err == 0)
     {
       /* Shift the outfiles array so the generated main comes first.
Index: gcc/gcc.c
===================================================================
--- gcc/gcc.c	(revision 227194)
+++ gcc/gcc.c	(working copy)
@@ -284,7 +284,7 @@ static const char *const spec_version = DEFAULT_TA
 static const char *spec_machine = DEFAULT_TARGET_MACHINE;
 static const char *spec_host_machine = DEFAULT_REAL_TARGET_MACHINE;
 
-/* List of offload targets.  */
+/* List of offload targets.  Empty string for -foffload=disable.  */
 
 static char *offload_targets = NULL;
 
@@ -400,6 +400,8 @@ static const char *compare_debug_auxbase_opt_spec_
 static const char *pass_through_libs_spec_func (int, const char **);
 static const char *replace_extension_spec_func (int, const char **);
 static const char *greater_than_spec_func (int, const char **);
+static const char *add_omp_infile_spec_func (int, const char **);
+
 static char *convert_white_space (char *);
 \f
 /* The Specs Language
@@ -1186,6 +1188,11 @@ static const char *const multilib_defaults_raw[] =
 
 static const char *const driver_self_specs[] = {
   "%{fdump-final-insns:-fdump-final-insns=.} %<fdump-final-insns",
+#ifdef ENABLE_OFFLOADING
+  /* If linking against libgomp, add a setup file.  */
+  "%{fopenacc|fopenmp|%:gt(%{ftree-parallelize-loops=*} 1):" \
+  "%:add-omp-infile()}",
+#endif /* ENABLE_OFFLOADING */
   DRIVER_SELF_SPECS, CONFIGURE_SPECS, GOMP_SELF_SPECS, GTM_SELF_SPECS,
   CILK_SELF_SPECS
 };
@@ -1613,6 +1620,7 @@ static const struct spec_function static_spec_func
   { "pass-through-libs",	pass_through_libs_spec_func },
   { "replace-extension",	replace_extension_spec_func },
   { "gt",			greater_than_spec_func },
+  { "add-omp-infile",		add_omp_infile_spec_func },
 #ifdef EXTRA_SPEC_FUNCTIONS
   EXTRA_SPEC_FUNCTIONS
 #endif
@@ -3209,7 +3217,8 @@ execute (void)
    The `validated' field describes whether any spec has looked at this switch;
    if it remains false at the end of the run, the switch must be meaningless.
    The `ordering' field is used to temporarily mark switches that have to be
-   kept in a specific order.  */
+   kept in a specific order.
+   The `lang_mask' field stores the flags associated with this option.  */
 
 #define SWITCH_LIVE    			(1 << 0)
 #define SWITCH_FALSE   			(1 << 1)
@@ -3225,6 +3234,7 @@ struct switchstr
   bool known;
   bool validated;
   bool ordering;
+  unsigned int lang_mask;
 };
 
 static struct switchstr *switches;
@@ -3233,6 +3243,10 @@ static int n_switches;
 
 static int n_switches_alloc;
 
+/* If nonzero, do not pass through switches for languages not matching
+   this mask.  */
+static unsigned int spec_lang_mask_accept;
+
 /* Set to zero if -fcompare-debug is disabled, positive if it's
    enabled and we're running the first compilation, negative if it's
    enabled and we're running the second compilation.  For most of the
@@ -3270,6 +3284,7 @@ struct infile
   const char *name;
   const char *language;
   struct compiler *incompiler;
+  unsigned int lang_mask;
   bool compiled;
   bool preprocessed;
 };
@@ -3463,15 +3478,16 @@ alloc_infile (void)
     }
 }
 
-/* Store an input file with the given NAME and LANGUAGE in
+/* Store an input file with the given NAME and LANGUAGE and LANG_MASK in
    infiles.  */
 
 static void
-add_infile (const char *name, const char *language)
+add_infile (const char *name, const char *language, unsigned int lang_mask)
 {
   alloc_infile ();
   infiles[n_infiles].name = name;
-  infiles[n_infiles++].language = language;
+  infiles[n_infiles].language = language;
+  infiles[n_infiles++].lang_mask = lang_mask;
 }
 
 /* Allocate space for a switch in switches.  */
@@ -3492,11 +3508,12 @@ alloc_switch (void)
 }
 
 /* Save an option OPT with N_ARGS arguments in array ARGS, marking it
-   as validated if VALIDATED and KNOWN if it is an internal switch.  */
+   as validated if VALIDATED and KNOWN if it is an internal switch.
+   LANG_MASK is the flags associated with this option.  */
 
 static void
 save_switch (const char *opt, size_t n_args, const char *const *args,
-	     bool validated, bool known)
+	     bool validated, bool known, unsigned int lang_mask)
 {
   alloc_switch ();
   switches[n_switches].part1 = opt + 1;
@@ -3513,6 +3530,7 @@ save_switch (const char *opt, size_t n_args, const
   switches[n_switches].validated = validated;
   switches[n_switches].known = known;
   switches[n_switches].ordering = 0;
+  switches[n_switches].lang_mask = lang_mask;
   n_switches++;
 }
 
@@ -3530,7 +3548,8 @@ driver_unknown_option_callback (const struct cl_de
 	 diagnosed only if there are warnings.  */
       save_switch (decoded->canonical_option[0],
 		   decoded->canonical_option_num_elements - 1,
-		   &decoded->canonical_option[1], false, true);
+		   &decoded->canonical_option[1], false, true,
+		   cl_options[decoded->opt_index].flags);
       return false;
     }
   if (decoded->opt_index == OPT_SPECIAL_unknown)
@@ -3538,7 +3557,8 @@ driver_unknown_option_callback (const struct cl_de
       /* Give it a chance to define it a spec file.  */
       save_switch (decoded->canonical_option[0],
 		   decoded->canonical_option_num_elements - 1,
-		   &decoded->canonical_option[1], false, false);
+		   &decoded->canonical_option[1], false, false,
+		   cl_options[decoded->opt_index].flags);
       return false;
     }
   else
@@ -3565,7 +3585,8 @@ driver_wrong_lang_callback (const struct cl_decode
   else
     save_switch (decoded->canonical_option[0],
 		 decoded->canonical_option_num_elements - 1,
-		 &decoded->canonical_option[1], false, true);
+		 &decoded->canonical_option[1], false, true,
+		 option->flags);
 }
 
 static const char *spec_lang = 0;
@@ -3815,7 +3836,8 @@ driver_handle_option (struct gcc_options *opts,
 	compare_debug_opt = NULL;
       else
 	compare_debug_opt = arg;
-      save_switch (compare_debug_replacement_opt, 0, NULL, validated, true);
+      save_switch (compare_debug_replacement_opt, 0, NULL, validated, true,
+		   cl_options[opt_index].flags);
       return true;
 
     case OPT_fdiagnostics_color_:
@@ -3870,17 +3892,17 @@ driver_handle_option (struct gcc_options *opts,
 	for (j = 0; arg[j]; j++)
 	  if (arg[j] == ',')
 	    {
-	      add_infile (save_string (arg + prev, j - prev), "*");
+	      add_infile (save_string (arg + prev, j - prev), "*", 0);
 	      prev = j + 1;
 	    }
 	/* Record the part after the last comma.  */
-	add_infile (arg + prev, "*");
+	add_infile (arg + prev, "*", 0);
       }
       do_save = false;
       break;
 
     case OPT_Xlinker:
-      add_infile (arg, "*");
+      add_infile (arg, "*", 0);
       do_save = false;
       break;
 
@@ -3897,19 +3919,21 @@ driver_handle_option (struct gcc_options *opts,
     case OPT_l:
       /* POSIX allows separation of -l and the lib arg; canonicalize
 	 by concatenating -l with its arg */
-      add_infile (concat ("-l", arg, NULL), "*");
+      add_infile (concat ("-l", arg, NULL), "*", 0);
       do_save = false;
       break;
 
     case OPT_L:
       /* Similarly, canonicalize -L for linkers that may not accept
 	 separate arguments.  */
-      save_switch (concat ("-L", arg, NULL), 0, NULL, validated, true);
+      save_switch (concat ("-L", arg, NULL), 0, NULL, validated, true,
+		   cl_options[opt_index].flags);
       return true;
 
     case OPT_F:
       /* Likewise -F.  */
-      save_switch (concat ("-F", arg, NULL), 0, NULL, validated, true);
+      save_switch (concat ("-F", arg, NULL), 0, NULL, validated, true,
+		   cl_options[opt_index].flags);
       return true;
 
     case OPT_save_temps:
@@ -4032,7 +4056,8 @@ driver_handle_option (struct gcc_options *opts,
       save_temps_prefix = xstrdup (arg);
       /* On some systems, ld cannot handle "-o" without a space.  So
 	 split the option from its argument.  */
-      save_switch ("-o", 1, &arg, validated, true);
+      save_switch ("-o", 1, &arg, validated, true,
+		   cl_options[opt_index].flags);
       return true;
 
 #ifdef ENABLE_DEFAULT_PIE
@@ -4068,7 +4093,8 @@ driver_handle_option (struct gcc_options *opts,
   if (do_save)
     save_switch (decoded->canonical_option[0],
 		 decoded->canonical_option_num_elements - 1,
-		 &decoded->canonical_option[1], validated, true);
+		 &decoded->canonical_option[1], validated, true,
+		 cl_options[opt_index].flags);
   return true;
 }
 
@@ -4365,7 +4391,7 @@ process_command (unsigned int decoded_options_coun
           if (strcmp (fname, "-") != 0 && access (fname, F_OK) < 0)
 	    perror_with_name (fname);
           else
-	    add_infile (arg, spec_lang);
+	    add_infile (arg, spec_lang, 0);
 
           free (fname);
 	  continue;
@@ -4376,6 +4402,11 @@ process_command (unsigned int decoded_options_coun
 			   CL_DRIVER, &handlers, global_dc);
     }
 
+  /* If the user didn't specify any, default to all configured offload
+     targets.  */
+  if (offload_targets == NULL)
+    handle_foffload_option (OFFLOAD_TARGETS);
+
   if (output_file
       && strcmp (output_file, "-") != 0
       && strcmp (output_file, HOST_BIT_BUCKET) != 0)
@@ -4507,7 +4538,8 @@ process_command (unsigned int decoded_options_coun
   if (compare_debug == 2 || compare_debug == 3)
     {
       const char *opt = concat ("-fcompare-debug=", compare_debug_opt, NULL);
-      save_switch (opt, 0, NULL, false, true);
+      save_switch (opt, 0, NULL, false, true,
+		   cl_options[OPT_fcompare_debug_].flags);
       compare_debug = 1;
     }
 
@@ -4518,7 +4550,7 @@ process_command (unsigned int decoded_options_coun
 
       /* Create a dummy input file, so that we can pass
 	 the help option on to the various sub-processes.  */
-      add_infile ("help-dummy", "c");
+      add_infile ("help-dummy", "c", 0);
     }
 
   alloc_switch ();
@@ -4719,13 +4751,15 @@ insert_wrapper (const char *wrapper)
 }
 
 /* Process the spec SPEC and run the commands specified therein.
+   If LANG_MASK is nonzero, switches for other languages are discarded.
    Returns 0 if the spec is successfully processed; -1 if failed.  */
 
 int
-do_spec (const char *spec)
+do_spec (const char *spec, unsigned int lang_mask)
 {
   int value;
 
+  spec_lang_mask_accept = lang_mask;
   value = do_spec_2 (spec);
 
   /* Force out any unfinished command.
@@ -4883,7 +4917,8 @@ do_self_spec (const char *spec)
 	      save_switch (decoded_options[j].canonical_option[0],
 			   (decoded_options[j].canonical_option_num_elements
 			    - 1),
-			   &decoded_options[j].canonical_option[1], false, true);
+			   &decoded_options[j].canonical_option[1], false, true,
+			   cl_options[decoded_options[j].opt_index].flags);
 	      break;
 
 	    default:
@@ -6479,6 +6514,14 @@ check_live_switch (int switchnum, int prefix_lengt
 static void
 give_switch (int switchnum, int omit_first_word)
 {
+  int lang_mask = switches[switchnum].lang_mask & ((1U << cl_lang_count) - 1);
+  unsigned int lang_mask_accept = (1U << cl_lang_count) - 1;
+  if (spec_lang_mask_accept != 0)
+    lang_mask_accept = spec_lang_mask_accept;
+  /* Drop switches specific to a language not in the given mask.  */
+  if (lang_mask != 0 && !(lang_mask & lang_mask_accept))
+    return;
+
   if ((switches[switchnum].live_cond & SWITCH_IGNORE) != 0)
     return;
 
@@ -7572,22 +7615,14 @@ driver::maybe_putenv_COLLECT_LTO_WRAPPER () const
 void
 driver::maybe_putenv_OFFLOAD_TARGETS () const
 {
-  const char *targets = offload_targets;
-
-  /* If no targets specified by -foffload, use all available targets.  */
-  if (!targets)
-    targets = OFFLOAD_TARGETS;
-
-  if (strlen (targets) > 0)
+  if (offload_targets && offload_targets[0] != '\0')
     {
       obstack_grow (&collect_obstack, "OFFLOAD_TARGET_NAMES=",
 		    sizeof ("OFFLOAD_TARGET_NAMES=") - 1);
-      obstack_grow (&collect_obstack, targets,
-		    strlen (targets) + 1);
+      obstack_grow (&collect_obstack, offload_targets,
+		    strlen (offload_targets) + 1);
       xputenv (XOBFINISH (&collect_obstack, char *));
     }
-
-  free (offload_targets);
 }
 
 /* Reject switches that no pass was interested in.  */
@@ -7891,7 +7926,8 @@ driver::do_spec_on_infiles () const
 		  debug_check_temp_file[1] = NULL;
 		}
 
-	      value = do_spec (input_file_compiler->spec);
+	      value = do_spec (input_file_compiler->spec,
+			       infiles[i].lang_mask);
 	      infiles[i].compiled = true;
 	      if (value < 0)
 		this_file_error = 1;
@@ -7905,7 +7941,8 @@ driver::do_spec_on_infiles () const
 		  n_switches_alloc = n_switches_alloc_debug_check[1];
 		  switches = switches_debug_check[1];
 
-		  value = do_spec (input_file_compiler->spec);
+		  value = do_spec (input_file_compiler->spec,
+				   infiles[i].lang_mask);
 
 		  compare_debug = -compare_debug;
 		  n_switches = n_switches_debug_check[0];
@@ -8060,7 +8097,7 @@ driver::maybe_run_linker (const char *argv0) const
 		    " to the linker.\n\n"));
 	  fflush (stdout);
 	}
-      int value = do_spec (link_command_spec);
+      int value = do_spec (link_command_spec, 0);
       if (value < 0)
 	errorcount = 1;
       linker_was_run = (tmp != execution_count);
@@ -9651,6 +9688,50 @@ greater_than_spec_func (int argc, const char **arg
   return NULL;
 }
 
+/* If applicable, generate a C source file containing a constructor call to
+   GOMP_set_offload_targets, to inform libgomp which offload targets have
+   actually been requested (-foffload=[...]), and adds that as an infile.  */
+
+static const char *
+add_omp_infile_spec_func (int argc, const char **)
+{
+  gcc_assert (argc == 0);
+  gcc_assert (offload_targets != NULL);
+
+  /* Nothing to do if we're not actually linking.  */
+  if (have_c)
+    return NULL;
+
+  int err;
+  const char *tmp_filename;
+  tmp_filename = make_temp_file (".c");
+  record_temp_file (tmp_filename, !save_temps_flag, 0);
+  FILE *f = fopen (tmp_filename, "w");
+  if (f == NULL)
+    fatal_error (input_location,
+		 "could not open temporary file %s", tmp_filename);
+  /* As libgomp uses constructors internally, and this code is only added when
+     linking against libgomp, it is fine to use a constructor here.  */
+  err = fprintf (f,
+		 "extern void GOMP_set_offload_targets (const char *);\n"
+		 "static __attribute__ ((constructor)) void\n"
+		 "init (void)\n"
+		 "{\n"
+		 "  GOMP_set_offload_targets (\"%s\");\n"
+		 "}\n",
+		 offload_targets);
+  if (err < 0)
+    fatal_error (input_location,
+		 "could not write to temporary file %s", tmp_filename);
+  err = fclose (f);
+  if (err == EOF)
+    fatal_error (input_location,
+		 "could not close temporary file %s", tmp_filename);
+
+  add_infile (tmp_filename, "cpp-output", CL_C);
+  return NULL;
+}
+
 /* Insert backslash before spaces in ORIG (usually a file path), to 
    avoid being broken by spec parser.
 
Index: gcc/gcc.h
===================================================================
--- gcc/gcc.h	(revision 227194)
+++ gcc/gcc.h	(working copy)
@@ -68,7 +68,7 @@ struct spec_function
 };
 
 /* These are exported by gcc.c.  */
-extern int do_spec (const char *);
+extern int do_spec (const char *, unsigned int);
 extern void record_temp_file (const char *, int, int);
 extern void pfatal_with_name (const char *) ATTRIBUTE_NORETURN;
 extern void set_input (const char *);
Index: gcc/fortran/gfortranspec.c
===================================================================
--- gcc/fortran/gfortranspec.c	(revision 227194)
+++ gcc/fortran/gfortranspec.c	(working copy)
@@ -439,7 +439,7 @@ int
 lang_specific_pre_link (void)
 {
   if (library)
-    do_spec ("%:include(libgfortran.spec)");
+    do_spec ("%:include(libgfortran.spec)", 0);
 
   return 0;
 }

-- 
Joseph S. Myers
joseph@codesourcery.com

^ permalink raw reply	[flat|nested] 62+ messages in thread
end of thread, other threads:[~2018-05-20 19:46 UTC | newest]

Thread overview: 62+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2014-09-27 18:17 [PATCH 2/n] OpenMP 4.0 offloading infrastructure: LTO streaming Ilya Verbin
2014-09-29  1:10 ` Jan Hubicka
2014-09-29 17:37   ` Ilya Verbin
2014-09-30 11:40     ` Thomas Schwinge
2014-10-01 16:13       ` Ilya Verbin
2014-10-08  8:45         ` Jakub Jelinek
2014-10-08  9:13           ` Jakub Jelinek
2014-10-15 14:28         ` Richard Biener
2014-10-20 11:21           ` Ilya Verbin
2014-10-20 11:26             ` Jakub Jelinek
2014-10-24 14:16             ` Ilya Verbin
2014-10-24 14:29               ` Jakub Jelinek
2014-10-28 19:32                 ` Ilya Verbin
2014-11-03  9:24                   ` Jakub Jelinek
2014-11-05 12:47                     ` Ilya Verbin
2014-11-05 12:50                       ` Jakub Jelinek
2014-11-07 14:41                         ` Kirill Yukhin
2014-11-12  9:32                       ` Richard Biener
2014-11-12 14:11                         ` Kirill Yukhin
2014-11-12 14:23                           ` Richard Biener
2014-11-12 14:35                             ` Kirill Yukhin
2014-11-12 14:41                               ` Richard Biener
2014-11-12 17:38                                 ` Ilya Verbin
2014-11-13  8:51                                   ` Richard Biener
2015-07-31 15:37                       ` Thomas Schwinge
2015-07-31 15:43                         ` Ilya Verbin
2015-08-05  8:40                           ` Richard Biener
2015-08-05 15:09                             ` Ilya Verbin
2015-08-14  9:49                               ` Forwarding -foffload=[...] from the driver (compile-time) to libgomp (run-time) (was: [PATCH 2/n] OpenMP 4.0 offloading infrastructure: LTO streaming) Thomas Schwinge
2015-08-14 13:29                                 ` Ilya Verbin
2015-08-17 13:57                                   ` Martin Jambor
2015-08-14 17:08                                 ` Joseph Myers
2015-08-14 21:48                                   ` Forwarding -foffload=[...] from the driver (compile-time) to libgomp (run-time) Thomas Schwinge
2015-08-15  4:03                                     ` Joseph Myers
2015-08-18 16:55                                       ` Thomas Schwinge
2015-08-20 23:38                                         ` Joseph Myers
2015-08-21 16:13                                           ` Nathan Sidwell
2015-08-21 16:21                                             ` Joseph Myers
2015-08-24 18:05                                               ` Joseph Myers
2015-08-24 22:50                                                 ` Joseph Myers
2015-08-24 23:26                                                   ` Nathan Sidwell
2015-08-25 15:04                                           ` Joseph Myers
2018-05-20 20:30                                           ` [og7] " Thomas Schwinge
2015-08-27 20:58 Pass -foffload targets from driver to libgomp at link time Joseph Myers
2015-09-03 14:58 ` Ping " Joseph Myers
2015-09-10 14:01   ` Ping^2 " Joseph Myers
2015-09-10 14:03     ` Bernd Schmidt
2015-09-11 14:29       ` Joseph Myers
2015-09-11 14:48         ` Bernd Schmidt
2015-09-11 15:28           ` Joseph Myers
2015-09-11 15:47             ` Jakub Jelinek
2015-09-11 16:16               ` Joseph Myers
2015-09-28 10:09               ` Thomas Schwinge
2015-09-29  9:48                 ` Jakub Jelinek
2015-09-30 16:15                   ` Thomas Schwinge
2015-10-19 16:56                     ` Forwarding -foffload=[...] from the driver (compile-time) to libgomp (run-time) Thomas Schwinge
2015-10-20 10:03                       ` Jakub Jelinek
2015-10-20 10:44                         ` Bernd Schmidt
2015-10-20 11:18                         ` Thomas Schwinge
2015-10-20 11:49                           ` Bernd Schmidt
2015-10-20 12:13                             ` Jakub Jelinek
2015-10-20 11:52                           ` Jakub Jelinek

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).