public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [gomp4] Preserve NVPTX "reconvergence" points
@ 2015-05-28 14:20 Julian Brown
  2015-05-28 14:59 ` Jakub Jelinek
  2015-05-28 15:02 ` Richard Biener
  0 siblings, 2 replies; 23+ messages in thread
From: Julian Brown @ 2015-05-28 14:20 UTC (permalink / raw)
  To: gcc-patches, Bernd Schmidt, Jakub Jelinek, Thomas Schwinge

[-- Attachment #1: Type: text/plain, Size: 6991 bytes --]

For NVPTX, it is vitally important that the divergence of threads
within a warp can be controlled: in particular we must be able to
generate code that we know "reconverges" at a particular point.
Unfortunately GCC's middle-end optimisers can cause this property to
be violated, which causes problems for the OpenACC execution model
we're planning to use for NVPTX.

As a brief example: code running in vector-single mode runs on a
single thread of a warp, and must broadcast condition results to other
threads of the warp so that they can "follow along" and be ready for
vector-partitioned execution when necessary.

#pragma acc parallel
{
  #pragma acc loop gang
  for (i = 0; i < N; i++)
  {
    /* This is vector-single mode.  */
    n = ...;
    switch (n)
    {
    case 1:
      #pragma acc loop vector
      for (...)
      {
        /* This is vector-partitioned mode.  */
      }
      ...
    }
  }
}

Here, the calculation "n = ..." takes place on a single thread (of
each partitioned gang of the outer loop), but the switch statement
(terminating the BB) must be executed by all threads in the warp. The
vector-single statements will be translated using a branch around for
the "idle" threads:

if (threadIdx.x == 0)
{
  n_0 = ...;
}
n_x = broadcast (n_0)
switch (n_x)
...

Where "broadcast" is an operation that transfers values from some
other thread of a warp (i.e., the zeroth) to the current thread
(implemented as a "shfl" instruction for NVPTX).

I observed a similar example to this cloning the broadcast and switch
instructions (in the .dom1 dump), along the lines of:

if (threadIdx.x == 0)
{
  n_0 = ...;
  n_x = broadcast (n_0)
  switch (n_x)
  ...
}
else
{
  n_x = broadcast (n_0)
  switch (n_x)
  ...
}

This doesn't work because the "broadcast" operation has to be run with
non-diverged warps for correct operation, and here there is divergence
due to the "if (threadIdx.x == 0)" condition.

So, the way I have tried to handle this is by attempting to inhibit
optimisation along edges which have a reconvergence point as their
destination. The essential idea is to make such edges "abnormal",
although the existing EDGE_ABNORMAL flag is not used because that has
implicit meaning built into it already, and the new edge type may need
to be handled differently in some areas. One example is that at
present, blocks concluding with GIMPLE_COND cannot have EDGE_ABNORMAL
set on their EDGE_TRUE or EDGE_FALSE outgoing edges.

The attached patch introduces a new edge flag (EDGE_TO_RECONVERGENCE),
for the GIMPLE CFG only. In principle there's nothing to stop the flag
being propagated to the RTL CFG also, in which case it'd probably be
set at the same time as EDGE_ABNORMAL, mirroring the semantics of e.g.
EDGE_EH, EDGE_ABNORMAL_CALL and EDGE_SIBCALL. Then, passes which
inspect the RTL CFG can continue to only check the ABNORMAL flag. But
so far (in rather limited testing!), that has not been observed to be
necessary. (We can control RTL CFG manipulation indirectly by using the
CANNOT_COPY_INSN_P target hook, sensitive e.g. to the "broadcast"
instruction.)

For the GIMPLE CFG (i.e. in passes operating on GIMPLE form),
EDGE_TO_RECONVERGENCE behaves mostly the same as EDGE_ABNORMAL (i.e.,
inhibiting certain optimisations), and so has been added to relevant
conditionals largely mechanically. Places where it is treated specially
are:

* tree-cfg.c:gimple_verify_flow_info does not permit EDGE_ABNORMAL on
  outgoing edges of a block concluding with a GIMPLE_COND statement.
  But, we allow EDGE_TO_RECONVERGENCE there.

* tree-vrp.c:find_conditional_asserts skips over outgoing GIMPLE_COND
  edges with EDGE_TO_RECONVERGENCE set (avoiding an ICE when the pass
  tries to split the edge later).

There are probably other optimisations that will be tripped up by the
new flag along the same lines as the VRP tweak above, which we will no
doubt discover in due course.

Together with the patch,

  https://gcc.gnu.org/ml/gcc-patches/2015-05/msg02612.html

This shows no regressions for the libgomp tests.

OK for gomp4 branch?

Thanks,

Julian

    ChangeLog

    gcc/
    * basic-block.h (EDGE_COMPLEX): Add EDGE_TO_RECONVERGENCE flag.
    (bb_hash_abnorm_or_reconv_pred): New function.
    (hash_abnormal_or_eh_outgoing_edge_p): Consider
    EDGE_TO_RECONVERGENCE also.
    * cfg-flags.def (TO_RECONVERGENCE): Add flag.
    * omp-low.c (predicate_bb): Set EDGE_TO_RECONVERGENCE on edges
    leading to a reconvergence point.
    * cfgbuild.c (purge_dead_tablejump_edges): Consider
    EDGE_TO_RECONVERGENCE.
    * cfgcleanup.c (try_crossjump_to_edge, try_head_merge_bb): Likewise.
    * cfgexpand.c (expand_gimple_tailcall, construct_exit_block)
    (pass_expand::execute): Likewise.
    * cfghooks.c (can_copy_bbs_p): Likewise.
    * cfgloop.c (bb_loop_header_p): Likewise.
    * cfgloopmanip.c (scale_loop_profile): Likewise.
    * gimple-iterator.c (gimple_find_edge_insert_loc): Likewise.
    * graph.c (draw_cfg_node_succ_edges): Likewise.
    * graphite-scope-detection.c (canonicalize_loop_closed_ssa):
    Likewise.
    * predict.c (tree_bb_level_predictions): Likewise.
    * profile.c (instrument_edges, branch_prop, find_spanning_tree):
    Likewise.
    * tree-cfg.c (replace_uses_by, gimple_split_edge)
    (gimple_redirect_edge_and_branch, split_critical_edges): Likewise.
    * tree-cfgcleanup.c (tree_forwarder_block_p, remove_forwarder_block)
    (pass_merge_phi::execute): Likewise.
    * tree-chkp.c (chkp_fix_cfg): Likewise.
    * tree-if-conv.c (if_convertible_bb_p): Likewise.
    * tree-inline.c (update_ssa_across_abnormal_edges): Likewise.
    * tree-into-ssa.c (rewrite_update_phi_arguments)
    (rewrite_update_dom_walker::before_dom_children)
    (create_new_def_for): Likewise.
    * tree-outof-ssa.c (eliminate_phi): Likewise.
    * tree-phinodes.c (add_phi_arg): Likewise.
    * tree-ssa-coalesce (coalesce_cost_edge, create_outofssa_var_map)
    (coalesce_partitions): Likewise.
    * tree-ssa-dom.c (cprop_into_successor_phis)
    (dom_opt_dom_walker::after_dom_children, propagate_rhs_into_lhs):
    Likewise.
    * tree-ssa-loop-im.c (loop_suitable_for_sm): Likewise.
    * tree-ssa-loop-prefetch.c (emit_mfence_after_loop)
    (may_use_storent_in_loop_p): Likewise.
    * tree-ssa-phiopt.c (tree_ssa_phiopt_worker): Likewise.
    * tree-ssa-pre.c (compute_antic, insert_into_preds_of_block):
    Likewise.
    * tree-ssa-propagate.c (simulate_block, replace_phi_args_in):
    Likewise.
    * tree-ssa-sink.c (sink_code_in_bb): Likewise.
    * tree-ssa-threadedge.c (thread_across_edge): Likewise.
    * tree-ssa-threadupdate.c (thread_single_edge): Likewise.
    * tree-ssa-uninit.c (compute_control_dep_chain): Likewise.
    * tree-ssa.c (verify_phi_args): Likewise.
    * tree-vect-loop.c (vect_analyze_loop_form): Likewise.
    * value-prof.c (gimple_ic): Likewise.
    * tree-vrp.c (infer_value_range, process_assert_insertions_for):
    Likewise.
    (find_conditional_asserts): Skip over EDGE_TO_RECONVERGENCE edges.

[-- Attachment #2: to-reconvergence-4.diff --]
[-- Type: text/x-patch, Size: 31591 bytes --]

commit 472bd543b30356f7a4c59efc961f9f61b11ca197
Author: Julian Brown <julian@codesourcery.com>
Date:   Wed May 20 11:35:45 2015 -0700

    Introduce EDGE_TO_RECONVERGENCE, and tweak some uses of EDGE_ABNORMAL.

diff --git a/gcc/basic-block.h b/gcc/basic-block.h
index f28fa57..7fe25f0 100644
--- a/gcc/basic-block.h
+++ b/gcc/basic-block.h
@@ -70,7 +70,8 @@ enum cfg_edge_flags {
    Test the edge flags on EDGE_COMPLEX to detect all forms of "strange"
    control flow transfers.  */
 #define EDGE_COMPLEX \
-  (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL | EDGE_EH | EDGE_PRESERVE)
+  (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL | EDGE_EH | EDGE_PRESERVE \
+   | EDGE_TO_RECONVERGENCE)
 
 struct GTY(()) rtl_bb_info {
   /* The first insn of the block is embedded into bb->il.x.  */
@@ -559,6 +560,20 @@ bb_has_abnormal_pred (basic_block bb)
   return false;
 }
 
+static inline bool
+bb_has_abnorm_or_reconv_pred (basic_block bb)
+{
+  edge e;
+  edge_iterator ei;
+
+  FOR_EACH_EDGE (e, ei, bb->preds)
+    {
+      if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
+	return true;
+    }
+  return false;
+}
+
 /* Return the fallthru edge in EDGES if it exists, NULL otherwise.  */
 static inline edge
 find_fallthru_edge (vec<edge, va_gc> *edges)
@@ -629,9 +644,10 @@ has_abnormal_or_eh_outgoing_edge_p (basic_block bb)
   edge_iterator ei;
 
   FOR_EACH_EDGE (e, ei, bb->succs)
-    if (e->flags & (EDGE_ABNORMAL | EDGE_EH))
+    if (e->flags & (EDGE_ABNORMAL | EDGE_EH | EDGE_TO_RECONVERGENCE))
       return true;
 
   return false;
 }
+
 #endif /* GCC_BASIC_BLOCK_H */
diff --git a/gcc/cfg-flags.def b/gcc/cfg-flags.def
index eedcd69..fd51e2f 100644
--- a/gcc/cfg-flags.def
+++ b/gcc/cfg-flags.def
@@ -177,6 +177,10 @@ DEF_EDGE_FLAG(TM_UNINSTRUMENTED, 15)
 /* Abort (over) edge out of a GIMPLE_TRANSACTION statement.  */
 DEF_EDGE_FLAG(TM_ABORT, 16)
 
+/* An "immutable" edge to an OpenACC (currently, NVPTX) reconvergence point. 
+   This flag is only used for the GIMPLE CFG.  */
+DEF_EDGE_FLAG(TO_RECONVERGENCE, 17)
+
 #endif
 
 /*
diff --git a/gcc/cfgbuild.c b/gcc/cfgbuild.c
index 7cbed50..7185f07 100644
--- a/gcc/cfgbuild.c
+++ b/gcc/cfgbuild.c
@@ -449,7 +449,7 @@ purge_dead_tablejump_edges (basic_block bb, rtx_jump_table_data *table)
       if (FULL_STATE (e->dest) & BLOCK_USED_BY_TABLEJUMP)
 	SET_STATE (e->dest, FULL_STATE (e->dest)
 			    & ~(size_t) BLOCK_USED_BY_TABLEJUMP);
-      else if (!(e->flags & (EDGE_ABNORMAL | EDGE_EH)))
+      else if (!(e->flags & (EDGE_ABNORMAL | EDGE_EH | EDGE_TO_RECONVERGENCE)))
 	{
 	  remove_edge (e);
 	  continue;
diff --git a/gcc/cfgcleanup.c b/gcc/cfgcleanup.c
index 797d14a..e73062a 100644
--- a/gcc/cfgcleanup.c
+++ b/gcc/cfgcleanup.c
@@ -2031,7 +2031,7 @@ try_crossjump_to_edge (int mode, edge e1, edge e2,
 
   /* Avoid deleting preserve label when redirecting ABNORMAL edges.  */
   if (block_has_preserve_label (e1->dest)
-      && (e1->flags & EDGE_ABNORMAL))
+      && (e1->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
     return false;
 
   /* Here we know that the insns in the end of SRC1 which are common with SRC2
@@ -2389,7 +2389,7 @@ try_head_merge_bb (basic_block bb)
 	  return false;
 	}
 
-      if (e->flags & EDGE_ABNORMAL)
+      if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	return false;
 
       /* Normally, all destination blocks must only be reachable from this
diff --git a/gcc/cfgexpand.c b/gcc/cfgexpand.c
index 5905ddb..688158b 100644
--- a/gcc/cfgexpand.c
+++ b/gcc/cfgexpand.c
@@ -3569,7 +3569,7 @@ expand_gimple_tailcall (basic_block bb, gcall *stmt, bool *can_fallthru)
 
   for (ei = ei_start (bb->succs); (e = ei_safe_edge (ei)); )
     {
-      if (!(e->flags & (EDGE_ABNORMAL | EDGE_EH)))
+      if (!(e->flags & (EDGE_ABNORMAL | EDGE_EH | EDGE_TO_RECONVERGENCE)))
 	{
 	  if (e->dest != EXIT_BLOCK_PTR_FOR_FN (cfun))
 	    {
@@ -5674,7 +5674,7 @@ construct_exit_block (void)
   while (ix < EDGE_COUNT (EXIT_BLOCK_PTR_FOR_FN (cfun)->preds))
     {
       e = EDGE_PRED (EXIT_BLOCK_PTR_FOR_FN (cfun), ix);
-      if (!(e->flags & EDGE_ABNORMAL))
+      if (!(e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
 	redirect_edge_succ (e, exit_block);
       else
 	ix++;
@@ -6222,7 +6222,7 @@ pass_expand::execute (function *fun)
 	     representation.  It is safe to remove them here as
 	     find_many_sub_basic_blocks will rediscover them.
 	     In the future we should get this fixed properly.  */
-	  if ((e->flags & EDGE_ABNORMAL)
+	  if ((e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	      && !(e->flags & EDGE_SIBCALL))
 	    remove_edge (e);
 	  else
diff --git a/gcc/cfghooks.c b/gcc/cfghooks.c
index fc23edb..64483f4 100644
--- a/gcc/cfghooks.c
+++ b/gcc/cfghooks.c
@@ -1304,7 +1304,7 @@ can_copy_bbs_p (basic_block *bbs, unsigned n)
       /* In case we should redirect abnormal edge during duplication, fail.  */
       edge_iterator ei;
       FOR_EACH_EDGE (e, ei, bbs[i]->succs)
-	if ((e->flags & EDGE_ABNORMAL)
+	if ((e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	    && (e->dest->flags & BB_DUPLICATED))
 	  {
 	    ret = false;
diff --git a/gcc/cfgloop.c b/gcc/cfgloop.c
index 5767494..d4a1c93 100644
--- a/gcc/cfgloop.c
+++ b/gcc/cfgloop.c
@@ -387,7 +387,7 @@ bb_loop_header_p (basic_block header)
 
   /* If we have an abnormal predecessor, do not consider the
      loop (not worth the problems).  */
-  if (bb_has_abnormal_pred (header))
+  if (bb_has_abnorm_or_reconv_pred (header))
     return false;
 
   /* Look for back edges where a predecessor is dominated
diff --git a/gcc/cfgloopmanip.c b/gcc/cfgloopmanip.c
index 45cc85d..8bbf4cb 100644
--- a/gcc/cfgloopmanip.c
+++ b/gcc/cfgloopmanip.c
@@ -554,7 +554,8 @@ scale_loop_profile (struct loop *loop, int scale, gcov_type iteration_bound)
 	  gcov_type count_delta;
 
           FOR_EACH_EDGE (other_e, ei, e->src->succs)
-	    if (!(other_e->flags & (EDGE_ABNORMAL | EDGE_FAKE))
+	    if (!(other_e->flags & (EDGE_ABNORMAL | EDGE_FAKE
+				    | EDGE_TO_RECONVERGENCE))
 		&& e != other_e)
 	      break;
 
diff --git a/gcc/gimple-iterator.c b/gcc/gimple-iterator.c
index df29123..bcb6649 100644
--- a/gcc/gimple-iterator.c
+++ b/gcc/gimple-iterator.c
@@ -792,7 +792,7 @@ gimple_find_edge_insert_loc (edge e, gimple_stmt_iterator *gsi,
      the last statement does not end a basic block, insert there.
      Except for the entry block.  */
   src = e->src;
-  if ((e->flags & EDGE_ABNORMAL) == 0
+  if ((e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)) == 0
       && single_succ_p (src)
       && src != ENTRY_BLOCK_PTR_FOR_FN (cfun))
     {
diff --git a/gcc/graph.c b/gcc/graph.c
index 5fb0d78..e5d3646 100644
--- a/gcc/graph.c
+++ b/gcc/graph.c
@@ -142,7 +142,7 @@ draw_cfg_node_succ_edges (pretty_printer *pp, int funcdef_no, basic_block bb)
 	  weight = 100;
 	}
 
-      if (e->flags & EDGE_ABNORMAL)
+      if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	color = "red";
 
       pp_printf (pp,
diff --git a/gcc/graphite-scop-detection.c b/gcc/graphite-scop-detection.c
index 02e9e50..8f91f6e 100644
--- a/gcc/graphite-scop-detection.c
+++ b/gcc/graphite-scop-detection.c
@@ -1339,7 +1339,7 @@ canonicalize_loop_closed_ssa (loop_p loop)
   edge e = single_exit (loop);
   basic_block bb;
 
-  if (!e || e->flags & EDGE_ABNORMAL)
+  if (!e || e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
     return;
 
   bb = e->dest;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index f82247b..e28f5b4 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -10447,6 +10447,7 @@ predicate_bb (basic_block bb, struct omp_region *parent)
       generate_vector_broadcast (broadcast_cond, cond_var, gsi_asgn);
 
       edge e = split_block (bb, asgn);
+      e->flags = EDGE_TO_RECONVERGENCE;
       skip_dest_bb = e->dest;
 
       gimple_cond_set_condition (as_a <gcond *> (stmt), EQ_EXPR,
@@ -10465,6 +10466,7 @@ predicate_bb (basic_block bb, struct omp_region *parent)
       generate_vector_broadcast (new_var, var, gsi_asgn);
 
       edge e = split_block (bb, asgn);
+      e->flags = EDGE_TO_RECONVERGENCE;
       skip_dest_bb = e->dest;
 
       gimple_switch_set_index (sstmt, new_var);
@@ -10477,6 +10479,7 @@ predicate_bb (basic_block bb, struct omp_region *parent)
 	  && gimple_code (stmt) != GIMPLE_OMP_CONTINUE)
 	{
 	  edge e = single_succ_edge (bb);
+	  e->flags = EDGE_TO_RECONVERGENCE;
 	  skip_dest_bb = e->dest;
 	  if (gimple_code (stmt) == GIMPLE_OMP_RETURN)
 	    {
@@ -10490,6 +10493,7 @@ predicate_bb (basic_block bb, struct omp_region *parent)
 	  if (!split_stmt)
 	    return;
 	  edge e = split_block (bb, split_stmt);
+	  e->flags = EDGE_TO_RECONVERGENCE;
 	  skip_dest_bb = e->dest;
 	  if (gimple_code (stmt) == GIMPLE_OMP_CONTINUE)
 	    {
@@ -10508,6 +10512,7 @@ predicate_bb (basic_block bb, struct omp_region *parent)
   else if (single_succ_p (bb))
     {
       edge e = single_succ_edge (bb);
+      e->flags |= EDGE_TO_RECONVERGENCE;
       skip_dest_bb = e->dest;
       if (gimple_code (stmt) == GIMPLE_GOTO)
 	gsi_prev (&gsi);
@@ -10540,7 +10545,8 @@ predicate_bb (basic_block bb, struct omp_region *parent)
       gsi_insert_after (&tmp_gsi, cond_stmt, GSI_CONTINUE_LINKING);
 
       e2->flags = EDGE_TRUE_VALUE;
-      make_edge (cond_bb, skip_dest_bb, EDGE_FALSE_VALUE);
+      make_edge (cond_bb, skip_dest_bb,
+		 EDGE_FALSE_VALUE | EDGE_TO_RECONVERGENCE);
     }
 }
 
diff --git a/gcc/predict.c b/gcc/predict.c
index 67d5d20..9098bce 100644
--- a/gcc/predict.c
+++ b/gcc/predict.c
@@ -2188,7 +2188,8 @@ tree_bb_level_predictions (void)
   edge_iterator ei;
 
   FOR_EACH_EDGE (e, ei, EXIT_BLOCK_PTR_FOR_FN (cfun)->preds)
-    if (!(e->flags & (EDGE_ABNORMAL | EDGE_FAKE | EDGE_EH)))
+    if (!(e->flags & (EDGE_ABNORMAL | EDGE_FAKE | EDGE_EH
+		      | EDGE_TO_RECONVERGENCE)))
       {
         has_return_edges = true;
 	break;
diff --git a/gcc/profile.c b/gcc/profile.c
index a178a1b..cd96b9f 100644
--- a/gcc/profile.c
+++ b/gcc/profile.c
@@ -170,7 +170,8 @@ instrument_edges (struct edge_list *el)
 
 	  if (!inf->ignore && !inf->on_tree)
 	    {
-	      gcc_assert (!(e->flags & EDGE_ABNORMAL));
+	      gcc_assert (!(e->flags & (EDGE_ABNORMAL
+					| EDGE_TO_RECONVERGENCE)));
 	      if (dump_file)
 		fprintf (dump_file, "Edge %d to %d instrumented%s\n",
 			 e->src->index, e->dest->index,
@@ -1107,7 +1108,8 @@ branch_prob (void)
 	      edge ne = single_succ_edge (new_bb);
 	      ne->goto_locus = e->goto_locus;
 	    }
-	  if ((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL))
+	  if ((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL
+			   | EDGE_TO_RECONVERGENCE))
 	       && e->dest != EXIT_BLOCK_PTR_FOR_FN (cfun))
 	    need_exit_edge = 1;
 	  if (e->dest == EXIT_BLOCK_PTR_FOR_FN (cfun))
@@ -1115,7 +1117,8 @@ branch_prob (void)
 	}
       FOR_EACH_EDGE (e, ei, bb->preds)
 	{
-	  if ((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL))
+	  if ((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL
+			   | EDGE_TO_RECONVERGENCE))
 	       && e->src != ENTRY_BLOCK_PTR_FOR_FN (cfun))
 	    need_entry_edge = 1;
 	  if (e->src == ENTRY_BLOCK_PTR_FOR_FN (cfun))
@@ -1180,7 +1183,8 @@ branch_prob (void)
       e->count = 0;
 
       /* Mark edges we've replaced by fake edges above as ignored.  */
-      if ((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL))
+      if ((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL
+		       | EDGE_TO_RECONVERGENCE))
 	  && e->src != ENTRY_BLOCK_PTR_FOR_FN (cfun)
 	  && e->dest != EXIT_BLOCK_PTR_FOR_FN (cfun))
 	{
@@ -1430,7 +1434,8 @@ find_spanning_tree (struct edge_list *el)
   for (i = 0; i < num_edges; i++)
     {
       edge e = INDEX_EDGE (el, i);
-      if (((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL | EDGE_FAKE))
+      if (((e->flags & (EDGE_ABNORMAL | EDGE_ABNORMAL_CALL | EDGE_FAKE
+			| EDGE_TO_RECONVERGENCE))
 	   || e->dest == EXIT_BLOCK_PTR_FOR_FN (cfun))
 	  && !EDGE_INFO (e)->ignore
 	  && (find_group (e->src) != find_group (e->dest)))
diff --git a/gcc/tree-cfg.c b/gcc/tree-cfg.c
index 99b27c7..fa12678 100644
--- a/gcc/tree-cfg.c
+++ b/gcc/tree-cfg.c
@@ -1865,7 +1865,7 @@ replace_uses_by (tree name, tree val)
 	    {
 	      e = gimple_phi_arg_edge (as_a <gphi *> (stmt),
 				       PHI_ARG_INDEX_FROM_USE (use));
-	      if (e->flags & EDGE_ABNORMAL
+	      if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)
 		  && !SSA_NAME_OCCURS_IN_ABNORMAL_PHI (val))
 		{
 		  /* This can only occur for virtual operands, since
@@ -2776,7 +2776,7 @@ gimple_split_edge (edge edge_in)
   edge new_edge, e;
 
   /* Abnormal edges cannot be split.  */
-  gcc_assert (!(edge_in->flags & EDGE_ABNORMAL));
+  gcc_assert (!(edge_in->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)));
 
   dest = edge_in->dest;
 
@@ -5578,7 +5578,7 @@ gimple_redirect_edge_and_branch (edge e, basic_block dest)
   edge ret;
   gimple stmt;
 
-  if (e->flags & EDGE_ABNORMAL)
+  if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
     return NULL;
 
   if (e->dest == dest)
@@ -5724,7 +5724,7 @@ gimple_redirect_edge_and_branch (edge e, basic_block dest)
 static bool
 gimple_can_remove_branch_p (const_edge e)
 {
-  if (e->flags & (EDGE_ABNORMAL | EDGE_EH))
+  if (e->flags & (EDGE_ABNORMAL | EDGE_EH | EDGE_TO_RECONVERGENCE))
     return false;
 
   return true;
@@ -8288,7 +8288,8 @@ split_critical_edges (void)
     {
       FOR_EACH_EDGE (e, ei, bb->succs)
         {
-	  if (EDGE_CRITICAL_P (e) && !(e->flags & EDGE_ABNORMAL))
+	  if (EDGE_CRITICAL_P (e) && !(e->flags & (EDGE_ABNORMAL
+						   | EDGE_TO_RECONVERGENCE)))
 	    split_edge (e);
 	  /* PRE inserts statements to edges and expects that
 	     since split_critical_edges was done beforehand, committing edge
@@ -8301,7 +8302,7 @@ split_critical_edges (void)
 	            || !gimple_seq_empty_p (phi_nodes (e->dest))
 		    || e->dest == EXIT_BLOCK_PTR_FOR_FN (cfun))
 		   && e->src != ENTRY_BLOCK_PTR_FOR_FN (cfun)
-	           && !(e->flags & EDGE_ABNORMAL))
+	           && !(e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
 	    {
 	      gimple_stmt_iterator gsi;
 
diff --git a/gcc/tree-cfgcleanup.c b/gcc/tree-cfgcleanup.c
index 26258aa..e361833 100644
--- a/gcc/tree-cfgcleanup.c
+++ b/gcc/tree-cfgcleanup.c
@@ -307,7 +307,8 @@ tree_forwarder_block_p (basic_block bb, bool phi_wanted)
       /* Nor should this be an infinite loop.  */
       || single_succ (bb) == bb
       /* BB may not have an abnormal outgoing edge.  */
-      || (single_succ_edge (bb)->flags & EDGE_ABNORMAL))
+      || (single_succ_edge (bb)->flags & (EDGE_ABNORMAL
+					  | EDGE_TO_RECONVERGENCE)))
     return false;
 
   gcc_checking_assert (bb != ENTRY_BLOCK_PTR_FOR_FN (cfun));
@@ -451,8 +452,8 @@ remove_forwarder_block (basic_block bb)
 
      So if there is an abnormal edge to BB, proceed only if there is
      no abnormal edge to DEST and there are no phi nodes in DEST.  */
-  if (bb_has_abnormal_pred (bb)
-      && (bb_has_abnormal_pred (dest)
+  if (bb_has_abnorm_or_reconv_pred (bb)
+      && (bb_has_abnorm_or_reconv_pred (dest)
 	  || !gimple_seq_empty_p (phi_nodes (dest))))
     return false;
 
@@ -483,7 +484,7 @@ remove_forwarder_block (basic_block bb)
     {
       bitmap_set_bit (cfgcleanup_altered_bbs, e->src->index);
 
-      if (e->flags & EDGE_ABNORMAL)
+      if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	{
 	  /* If there is an abnormal edge, redirect it anyway, and
 	     move the labels to the new block to make it legal.  */
@@ -1019,7 +1020,7 @@ pass_merge_phi::execute (function *fun)
       if (gimple_seq_empty_p (phi_nodes (dest))
 	  /* We don't want to deal with a basic block with
 	     abnormal edges.  */
-	  || bb_has_abnormal_pred (bb))
+	  || bb_has_abnorm_or_reconv_pred (bb))
 	continue;
 
       if (!dominated_by_p (CDI_DOMINATORS, dest, bb))
diff --git a/gcc/tree-chkp.c b/gcc/tree-chkp.c
index 288470b..a916525 100644
--- a/gcc/tree-chkp.c
+++ b/gcc/tree-chkp.c
@@ -3994,7 +3994,7 @@ chkp_fix_cfg ()
 	    /* We cannot split abnormal edge.  Therefore we
 	       store its params, make it regular and then
 	       rebuild abnormal edge after split.  */
-	    if (fall->flags & EDGE_ABNORMAL)
+	    if (fall->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	      {
 		flags = fall->flags & ~EDGE_FALLTHRU;
 		dest = fall->dest;
diff --git a/gcc/tree-if-conv.c b/gcc/tree-if-conv.c
index 49ff458..7849934 100644
--- a/gcc/tree-if-conv.c
+++ b/gcc/tree-if-conv.c
@@ -1070,7 +1070,8 @@ if_convertible_bb_p (struct loop *loop, basic_block bb, basic_block exit_bb)
 
   /* Be less adventurous and handle only normal edges.  */
   FOR_EACH_EDGE (e, ei, bb->succs)
-    if (e->flags & (EDGE_EH | EDGE_ABNORMAL | EDGE_IRREDUCIBLE_LOOP))
+    if (e->flags & (EDGE_EH | EDGE_ABNORMAL | EDGE_IRREDUCIBLE_LOOP
+		    | EDGE_TO_RECONVERGENCE))
       {
 	if (dump_file && (dump_flags & TDF_DETAILS))
 	  fprintf (dump_file, "Difficult to handle edges\n");
diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c
index 71d75d9..72aa414 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -2182,8 +2182,10 @@ update_ssa_across_abnormal_edges (basic_block bb, basic_block ret_bb,
 
 	    re = find_edge (ret_bb, e->dest);
 	    gcc_checking_assert (re);
-	    gcc_assert ((re->flags & (EDGE_EH | EDGE_ABNORMAL))
-			== (e->flags & (EDGE_EH | EDGE_ABNORMAL)));
+	    gcc_assert ((re->flags & (EDGE_EH | EDGE_ABNORMAL
+				      | EDGE_TO_RECONVERGENCE))
+			== (e->flags & (EDGE_EH | EDGE_ABNORMAL
+					| EDGE_TO_RECONVERGENCE)));
 
 	    SET_USE (PHI_ARG_DEF_PTR_FROM_EDGE (phi, e),
 		     USE_FROM_PTR (PHI_ARG_DEF_PTR_FROM_EDGE (phi, re)));
diff --git a/gcc/tree-into-ssa.c b/gcc/tree-into-ssa.c
index e8b55c1..7e9cb57 100644
--- a/gcc/tree-into-ssa.c
+++ b/gcc/tree-into-ssa.c
@@ -2101,7 +2101,7 @@ rewrite_update_phi_arguments (basic_block bb)
 	    }
 
 
-	  if (e->flags & EDGE_ABNORMAL)
+	  if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	    SSA_NAME_OCCURS_IN_ABNORMAL_PHI (USE_FROM_PTR (arg_p)) = 1;
 	}
     }
@@ -2138,7 +2138,7 @@ rewrite_update_dom_walker::before_dom_children (basic_block bb)
 
   /* Mark the LHS if any of the arguments flows through an abnormal
      edge.  */
-  is_abnormal_phi = bb_has_abnormal_pred (bb);
+  is_abnormal_phi = bb_has_abnorm_or_reconv_pred (bb);
 
   /* If any of the PHI nodes is a replacement for a name in
      OLD_SSA_NAMES or it's one of the names in NEW_SSA_NAMES, then
@@ -2899,7 +2899,8 @@ create_new_def_for (tree old_name, gimple stmt, def_operand_p def)
       basic_block bb = gimple_bb (stmt);
 
       /* If needed, mark NEW_NAME as occurring in an abnormal PHI node. */
-      SSA_NAME_OCCURS_IN_ABNORMAL_PHI (new_name) = bb_has_abnormal_pred (bb);
+      SSA_NAME_OCCURS_IN_ABNORMAL_PHI (new_name)
+        = bb_has_abnorm_or_reconv_pred (bb);
     }
 
   add_new_name_mapping (new_name, old_name);
diff --git a/gcc/tree-outof-ssa.c b/gcc/tree-outof-ssa.c
index e23bc0b..fe24cf8 100644
--- a/gcc/tree-outof-ssa.c
+++ b/gcc/tree-outof-ssa.c
@@ -766,7 +766,7 @@ eliminate_phi (edge e, elim_graph g)
   gcc_assert (g->copy_locus.length () == 0);
 
   /* Abnormal edges already have everything coalesced.  */
-  if (e->flags & EDGE_ABNORMAL)
+  if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
     return;
 
   g->e = e;
diff --git a/gcc/tree-phinodes.c b/gcc/tree-phinodes.c
index d657907..05c383d 100644
--- a/gcc/tree-phinodes.c
+++ b/gcc/tree-phinodes.c
@@ -392,7 +392,7 @@ add_phi_arg (gphi *phi, tree def, edge e, source_location locus)
 
   /* Copy propagation needs to know what object occur in abnormal
      PHI nodes.  This is a convenient place to record such information.  */
-  if (e->flags & EDGE_ABNORMAL)
+  if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
     {
       SSA_NAME_OCCURS_IN_ABNORMAL_PHI (def) = 1;
       SSA_NAME_OCCURS_IN_ABNORMAL_PHI (PHI_RESULT (phi)) = 1;
diff --git a/gcc/tree-ssa-coalesce.c b/gcc/tree-ssa-coalesce.c
index eeac5a4..a72c86c 100644
--- a/gcc/tree-ssa-coalesce.c
+++ b/gcc/tree-ssa-coalesce.c
@@ -170,7 +170,7 @@ coalesce_cost_edge (edge e)
   /* Inserting copy on critical edge costs more than inserting it elsewhere.  */
   if (EDGE_CRITICAL_P (e))
     mult = 2;
-  if (e->flags & EDGE_ABNORMAL)
+  if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
     return MUST_COALESCE_COST;
   if (e->flags & EDGE_EH)
     {
@@ -975,11 +975,11 @@ create_outofssa_var_map (coalesce_list_p cl, bitmap used_in_copy)
 
 	      register_ssa_partition (map, arg);
 	      if (gimple_can_coalesce_p (arg, res)
-		  || (e->flags & EDGE_ABNORMAL))
+		  || (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
 		{
 		  saw_copy = true;
 		  bitmap_set_bit (used_in_copy, SSA_NAME_VERSION (arg));
-		  if ((e->flags & EDGE_ABNORMAL) == 0)
+		  if ((e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)) == 0)
 		    {
 		      int cost = coalesce_cost_edge (e);
 		      if (cost == 1 && has_single_use (arg))
@@ -1206,7 +1206,7 @@ coalesce_partitions (var_map map, ssa_conflicts_p graph, coalesce_list_p cl,
   FOR_EACH_BB_FN (bb, cfun)
     {
       FOR_EACH_EDGE (e, ei, bb->preds)
-	if (e->flags & EDGE_ABNORMAL)
+	if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	  {
 	    gphi_iterator gsi;
 	    for (gsi = gsi_start_phis (bb); !gsi_end_p (gsi);
diff --git a/gcc/tree-ssa-dom.c b/gcc/tree-ssa-dom.c
index 14f3e9e..9a27acf 100644
--- a/gcc/tree-ssa-dom.c
+++ b/gcc/tree-ssa-dom.c
@@ -1874,7 +1874,7 @@ cprop_into_successor_phis (basic_block bb)
 
       /* If this is an abnormal edge, then we do not want to copy propagate
 	 into the PHI alternative associated with this edge.  */
-      if (e->flags & EDGE_ABNORMAL)
+      if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	continue;
 
       gsi = gsi_start_phis (e->dest);
@@ -1983,7 +1983,8 @@ dom_opt_dom_walker::after_dom_children (basic_block bb)
      may be able to statically determine which of the outgoing edges
      will be traversed when the incoming edge from BB is traversed.  */
   if (single_succ_p (bb)
-      && (single_succ_edge (bb)->flags & EDGE_ABNORMAL) == 0
+      && (single_succ_edge (bb)->flags & (EDGE_ABNORMAL
+					  | EDGE_TO_RECONVERGENCE)) == 0
       && potentially_threadable_block (single_succ (bb)))
     {
       thread_across_edge (single_succ_edge (bb));
@@ -1991,8 +1992,10 @@ dom_opt_dom_walker::after_dom_children (basic_block bb)
   else if ((last = last_stmt (bb))
 	   && gimple_code (last) == GIMPLE_COND
 	   && EDGE_COUNT (bb->succs) == 2
-	   && (EDGE_SUCC (bb, 0)->flags & EDGE_ABNORMAL) == 0
-	   && (EDGE_SUCC (bb, 1)->flags & EDGE_ABNORMAL) == 0)
+	   && (EDGE_SUCC (bb, 0)->flags & (EDGE_ABNORMAL
+					   | EDGE_TO_RECONVERGENCE)) == 0
+	   && (EDGE_SUCC (bb, 1)->flags & (EDGE_ABNORMAL
+					   | EDGE_TO_RECONVERGENCE)) == 0)
     {
       edge true_edge, false_edge;
 
@@ -2957,7 +2960,7 @@ propagate_rhs_into_lhs (gimple stmt, tree lhs, tree rhs, bitmap interesting_name
 
 		  /* And fixup the flags on the single remaining edge.  */
 		  te->flags &= ~(EDGE_TRUE_VALUE | EDGE_FALSE_VALUE);
-		  te->flags &= ~EDGE_ABNORMAL;
+		  te->flags &= ~(EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE);
 		  te->flags |= EDGE_FALLTHRU;
 		  if (te->probability > REG_BR_PROB_BASE)
 		    te->probability = REG_BR_PROB_BASE;
diff --git a/gcc/tree-ssa-loop-im.c b/gcc/tree-ssa-loop-im.c
index 11fc699..cb61979 100644
--- a/gcc/tree-ssa-loop-im.c
+++ b/gcc/tree-ssa-loop-im.c
@@ -2313,7 +2313,7 @@ loop_suitable_for_sm (struct loop *loop ATTRIBUTE_UNUSED,
   edge ex;
 
   FOR_EACH_VEC_ELT (exits, i, ex)
-    if (ex->flags & (EDGE_ABNORMAL | EDGE_EH))
+    if (ex->flags & (EDGE_ABNORMAL | EDGE_EH | EDGE_TO_RECONVERGENCE))
       return false;
 
   return true;
diff --git a/gcc/tree-ssa-loop-prefetch.c b/gcc/tree-ssa-loop-prefetch.c
index a948d68..f817b9e 100644
--- a/gcc/tree-ssa-loop-prefetch.c
+++ b/gcc/tree-ssa-loop-prefetch.c
@@ -1283,7 +1283,7 @@ emit_mfence_after_loop (struct loop *loop)
       if (!single_pred_p (exit->dest)
 	  /* If possible, we prefer not to insert the fence on other paths
 	     in cfg.  */
-	  && !(exit->flags & EDGE_ABNORMAL))
+	  && !(exit->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
 	split_loop_exit_edge (exit);
       bsi = gsi_after_labels (exit->dest);
 
@@ -1313,7 +1313,7 @@ may_use_storent_in_loop_p (struct loop *loop)
       edge exit;
 
       FOR_EACH_VEC_ELT (exits, i, exit)
-	if ((exit->flags & EDGE_ABNORMAL)
+	if ((exit->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	    && exit->dest == EXIT_BLOCK_PTR_FOR_FN (cfun))
 	  ret = false;
 
diff --git a/gcc/tree-ssa-phiopt.c b/gcc/tree-ssa-phiopt.c
index 7c846c2..91ae8dc 100644
--- a/gcc/tree-ssa-phiopt.c
+++ b/gcc/tree-ssa-phiopt.c
@@ -239,8 +239,8 @@ tree_ssa_phiopt_worker (bool do_store_elim, bool do_hoist_loads)
       bb2 = e2->dest;
 
       /* We cannot do the optimization on abnormal edges.  */
-      if ((e1->flags & EDGE_ABNORMAL) != 0
-          || (e2->flags & EDGE_ABNORMAL) != 0)
+      if ((e1->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)) != 0
+          || (e2->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)) != 0)
        continue;
 
       /* If either bb1's succ or bb2 or bb2's succ is non NULL.  */
diff --git a/gcc/tree-ssa-pre.c b/gcc/tree-ssa-pre.c
index d857d84..2f69389 100644
--- a/gcc/tree-ssa-pre.c
+++ b/gcc/tree-ssa-pre.c
@@ -2410,7 +2410,7 @@ compute_antic (void)
       FOR_EACH_EDGE (e, ei, block->preds)
 	{
 	  e->flags &= ~EDGE_DFS_BACK;
-	  if (e->flags & EDGE_ABNORMAL)
+	  if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	    {
 	      bitmap_set_bit (has_abnormal_preds, block->index);
 	      break;
@@ -3035,7 +3035,7 @@ insert_into_preds_of_block (basic_block block, unsigned int exprnum,
 	{
 	  builtexpr = create_expression_by_pieces (bprime, eprime,
 						   &stmts, type);
-	  gcc_assert (!(pred->flags & EDGE_ABNORMAL));
+	  gcc_assert (!(pred->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)));
 	  gsi_insert_seq_on_edge (pred, stmts);
 	  if (!builtexpr)
 	    {
diff --git a/gcc/tree-ssa-propagate.c b/gcc/tree-ssa-propagate.c
index e23da70..26b8ec8 100644
--- a/gcc/tree-ssa-propagate.c
+++ b/gcc/tree-ssa-propagate.c
@@ -526,7 +526,7 @@ simulate_block (basic_block block)
       normal_edge = NULL;
       FOR_EACH_EDGE (e, ei, block->succs)
 	{
-	  if (e->flags & (EDGE_ABNORMAL | EDGE_EH))
+	  if (e->flags & (EDGE_ABNORMAL | EDGE_EH | EDGE_TO_RECONVERGENCE))
 	    add_control_edge (e);
 	  else
 	    {
@@ -1072,7 +1072,7 @@ replace_phi_args_in (gphi *phi, ssa_prop_get_value_fn get_value)
 		 through an abnormal edge, update the replacement
 		 accordingly.  */
 	      if (TREE_CODE (val) == SSA_NAME
-		  && e->flags & EDGE_ABNORMAL
+		  && e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)
 		  && !SSA_NAME_OCCURS_IN_ABNORMAL_PHI (val))
 		{
 		  /* This can only occur for virtual operands, since
diff --git a/gcc/tree-ssa-sink.c b/gcc/tree-ssa-sink.c
index 1ed8a0e..c3d091c 100644
--- a/gcc/tree-ssa-sink.c
+++ b/gcc/tree-ssa-sink.c
@@ -501,7 +501,7 @@ sink_code_in_bb (basic_block bb)
 
   /* We can't move things across abnormal edges, so don't try.  */
   FOR_EACH_EDGE (e, ei, bb->succs)
-    if (e->flags & EDGE_ABNORMAL)
+    if (e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
       goto earlyout;
 
   for (gsi = gsi_last_bb (bb); !gsi_end_p (gsi);)
diff --git a/gcc/tree-ssa-threadedge.c b/gcc/tree-ssa-threadedge.c
index acbbb67..5b6626a 100644
--- a/gcc/tree-ssa-threadedge.c
+++ b/gcc/tree-ssa-threadedge.c
@@ -1417,7 +1417,7 @@ thread_across_edge (gcond *dummy_cond,
     /* If E->dest has abnormal outgoing edges, then there's no guarantee
        we can safely redirect any of the edges.  Just punt those cases.  */
     FOR_EACH_EDGE (taken_edge, ei, e->dest->succs)
-      if (taken_edge->flags & EDGE_ABNORMAL)
+      if (taken_edge->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
 	{
 	  const_and_copies->pop_to_marker ();
 	  BITMAP_FREE (visited);
diff --git a/gcc/tree-ssa-threadupdate.c b/gcc/tree-ssa-threadupdate.c
index 0d61c18..49b010d 100644
--- a/gcc/tree-ssa-threadupdate.c
+++ b/gcc/tree-ssa-threadupdate.c
@@ -1661,7 +1661,8 @@ thread_single_edge (edge e)
       remove_ctrl_stmt_and_useless_edges (bb, eto->dest);
 
       /* And fixup the flags on the single remaining edge.  */
-      eto->flags &= ~(EDGE_TRUE_VALUE | EDGE_FALSE_VALUE | EDGE_ABNORMAL);
+      eto->flags &= ~(EDGE_TRUE_VALUE | EDGE_FALSE_VALUE | EDGE_ABNORMAL
+		      | EDGE_TO_RECONVERGENCE);
       eto->flags |= EDGE_FALLTHRU;
 
       return bb;
diff --git a/gcc/tree-ssa-uninit.c b/gcc/tree-ssa-uninit.c
index 19a3e82..a23d83a 100644
--- a/gcc/tree-ssa-uninit.c
+++ b/gcc/tree-ssa-uninit.c
@@ -459,7 +459,7 @@ compute_control_dep_chain (basic_block bb, basic_block dep_bb,
     {
       basic_block cd_bb;
       int post_dom_check = 0;
-      if (e->flags & (EDGE_FAKE | EDGE_ABNORMAL))
+      if (e->flags & (EDGE_FAKE | EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
         continue;
 
       cd_bb = e->dest;
diff --git a/gcc/tree-ssa.c b/gcc/tree-ssa.c
index 10d3314..86fd30c 100644
--- a/gcc/tree-ssa.c
+++ b/gcc/tree-ssa.c
@@ -883,7 +883,9 @@ verify_phi_args (gphi *phi, basic_block bb, basic_block *definition_block)
 	{
 	  err = verify_ssa_name (op, virtual_operand_p (gimple_phi_result (phi)));
 	  err |= verify_use (e->src, definition_block[SSA_NAME_VERSION (op)],
-			     op_p, phi, e->flags & EDGE_ABNORMAL, NULL);
+			     op_p, phi,
+			     e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE),
+			     NULL);
 	}
 
       if (TREE_CODE (op) == ADDR_EXPR)
diff --git a/gcc/tree-vect-loop.c b/gcc/tree-vect-loop.c
index 49bf518..dcca3fe 100644
--- a/gcc/tree-vect-loop.c
+++ b/gcc/tree-vect-loop.c
@@ -1277,7 +1277,7 @@ vect_analyze_loop_form (struct loop *loop)
   if (!single_pred_p (single_exit (loop)->dest))
     {
       edge e = single_exit (loop);
-      if (!(e->flags & EDGE_ABNORMAL))
+      if (!(e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
 	{
 	  split_loop_exit_edge (e);
 	  if (dump_enabled_p ())
diff --git a/gcc/tree-vrp.c b/gcc/tree-vrp.c
index 6744a91..97a3ccf 100644
--- a/gcc/tree-vrp.c
+++ b/gcc/tree-vrp.c
@@ -4926,7 +4926,7 @@ infer_value_range (gimple stmt, tree op, enum tree_code *comp_code_p, tree *val_
       edge e;
 
       FOR_EACH_EDGE (e, ei, gimple_bb (stmt)->succs)
-	if (!(e->flags & EDGE_ABNORMAL))
+	if (!(e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
 	  break;
       if (e == NULL)
 	return false;
@@ -5949,7 +5949,9 @@ find_conditional_asserts (basic_block bb, gcond *last)
      insert.  */
   FOR_EACH_EDGE (e, ei, bb->succs)
     {
-      if (e->dest == bb)
+      /* Skip over EDGE_TO_RECONVERGENCE edges because they cannot be split
+	 later in the pass.  */
+      if (e->dest == bb || (e->flags & EDGE_TO_RECONVERGENCE))
 	continue;
 
       /* Register the necessary assertions for each operand in the
@@ -6412,7 +6414,7 @@ process_assert_insertions_for (tree name, assert_locus_t loc)
      STMT is not control flow, there may only be one non-abnormal edge
      out of BB.  */
   FOR_EACH_EDGE (e, ei, loc->bb->succs)
-    if (!(e->flags & EDGE_ABNORMAL))
+    if (!(e->flags & (EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE)))
       {
 	gsi_insert_on_edge (e, assert_stmt);
 	return true;
diff --git a/gcc/value-prof.c b/gcc/value-prof.c
index b16bce8..7a70c49 100644
--- a/gcc/value-prof.c
+++ b/gcc/value-prof.c
@@ -1565,7 +1565,7 @@ gimple_ic (gcall *icall_stmt, struct cgraph_node *direct_call,
     }
 
   FOR_EACH_EDGE (e_eh, ei, icall_bb->succs)
-    if (e_eh->flags & (EDGE_EH | EDGE_ABNORMAL))
+    if (e_eh->flags & (EDGE_EH | EDGE_ABNORMAL | EDGE_TO_RECONVERGENCE))
       {
 	e = make_edge (dcall_bb, e_eh->dest, e_eh->flags);
 	for (gphi_iterator psi = gsi_start_phis (e_eh->dest);

^ permalink raw reply	[flat|nested] 23+ messages in thread
end of thread, other threads:[~2015-06-24 13:53 UTC | newest]

Thread overview: 23+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-05-28 14:20 [gomp4] Preserve NVPTX "reconvergence" points Julian Brown
2015-05-28 14:59 ` Jakub Jelinek
2015-05-28 15:14   ` Thomas Schwinge
2015-05-28 15:28     ` Jakub Jelinek
2015-06-19 10:44       ` Bernd Schmidt
2015-06-19 12:32         ` Jakub Jelinek
2015-06-19 13:07           ` Bernd Schmidt
2015-06-19 14:10             ` Jakub Jelinek
2015-06-22 14:04               ` Bernd Schmidt
2015-06-22 14:25                 ` Jakub Jelinek
2015-06-24 13:37               ` Bernd Schmidt
2015-06-24 14:08                 ` Jakub Jelinek
2015-06-22 14:00           ` Julian Brown
2015-06-22 14:36             ` Jakub Jelinek
2015-06-22 15:18               ` Julian Brown
2015-06-22 15:33               ` Bernd Schmidt
2015-06-22 16:13                 ` Nathan Sidwell
2015-06-22 16:27                   ` Jakub Jelinek
2015-06-22 16:35                     ` Nathan Sidwell
2015-06-22 17:54               ` Julian Brown
2015-06-22 18:48                 ` Jakub Jelinek
2015-05-28 15:02 ` Richard Biener
2015-06-03 11:47   ` Julian Brown

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