[gomp4] Preserve NVPTX "reconvergence" points

[Richard Biener]

On Thu, May 28, 2015 at 4:06 PM, Julian Brown <julian@codesourcery.com> wrote:
> 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?

Hmm, I don't think adding a new edge flag is good nor necessary.  It seems to
me that instead the broadcast operation should have abnormal control flow
and thus basic-blocks should be split either before or after it (so either
incoming or outgoing edge(s) should be abnormal).  I suppose splitting
before the broadcast would be best (thus handle it similar to setjmp ()).

Richard.

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