OpenMP offloading to NVPTX: middle-end changes

	* internal-fn.c (expand_GOMP_SIMT_LANE): New.
	(expand_GOMP_SIMT_VF): New.
	(expand_GOMP_SIMT_LAST_LANE): New.
	(expand_GOMP_SIMT_ORDERED_PRED): New.
	(expand_GOMP_SIMT_VOTE_ANY): New.
	(expand_GOMP_SIMT_XCHG_BFLY): New.
	(expand_GOMP_SIMT_XCHG_IDX): New.
	* internal-fn.def (GOMP_SIMT_LANE): New.
	(GOMP_SIMT_VF): New.
	(GOMP_SIMT_LAST_LANE): New.
	(GOMP_SIMT_ORDERED_PRED): New.
	(GOMP_SIMT_VOTE_ANY): New.
	(GOMP_SIMT_XCHG_BFLY): New.
	(GOMP_SIMT_XCHG_IDX): New.
	* omp-low.c (omp_maybe_offloaded_ctx): New, outlined from...
	(create_omp_child_function): ...here.  Set "omp target entrypoint"
	or "omp declare target" attribute based on is_gimple_omp_offloaded.
	(omp_max_simt_vf): New.  Use it...
	(omp_max_vf): ...here.
	(lower_rec_input_clauses): Add reduction lowering for SIMT execution.
	(lower_lastprivate_clauses): Likewise, for "lastprivate" lowering.
	(lower_omp_ordered): Likewise, for "ordered" lowering.
	(expand_omp_simd): Add SIMT transforms.
	(pass_data_lower_omp): Add PROP_gimple_lomp_dev.
	(execute_omp_device_lower): New.
	(pass_data_omp_device_lower): New.
	(pass_omp_device_lower): New pass.
	(make_pass_omp_device_lower): New.
	* passes.def (pass_omp_device_lower): Position new pass.
	* tree-pass.h (PROP_gimple_lomp_dev): Define.
	(make_pass_omp_device_lower): Declare.

From-SVN: r242710

1 files changed, 384 insertions, 64 deletions

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 7c58c033ded..6c52bff74ba 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2427,6 +2427,20 @@ cilk_for_check_loop_diff_type (tree type)
     }
 }
 
+/* Return true if CTX may belong to offloaded code: either if current function
+   is offloaded, or any enclosing context corresponds to a target region.  */
+
+static bool
+omp_maybe_offloaded_ctx (omp_context *ctx)
+{
+  if (cgraph_node::get (current_function_decl)->offloadable)
+    return true;
+  for (; ctx; ctx = ctx->outer)
+    if (is_gimple_omp_offloaded (ctx->stmt))
+      return true;
+  return false;
+}
+
 /* Build a decl for the omp child function.  It'll not contain a body
    yet, just the bare decl.  */
 
@@ -2475,28 +2489,24 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
   DECL_CONTEXT (decl) = NULL_TREE;
   DECL_INITIAL (decl) = make_node (BLOCK);
   BLOCK_SUPERCONTEXT (DECL_INITIAL (decl)) = decl;
-  if (cgraph_node::get (current_function_decl)->offloadable)
-    cgraph_node::get_create (decl)->offloadable = 1;
-  else
+  if (omp_maybe_offloaded_ctx (ctx))
     {
-      omp_context *octx;
-      for (octx = ctx; octx; octx = octx->outer)
-	if (is_gimple_omp_offloaded (octx->stmt))
-	  {
-	    cgraph_node::get_create (decl)->offloadable = 1;
-	    if (ENABLE_OFFLOADING)
-	      g->have_offload = true;
-
-	    break;
-	  }
+      cgraph_node::get_create (decl)->offloadable = 1;
+      if (ENABLE_OFFLOADING)
+	g->have_offload = true;
     }
 
   if (cgraph_node::get_create (decl)->offloadable
       && !lookup_attribute ("omp declare target",
                            DECL_ATTRIBUTES (current_function_decl)))
-    DECL_ATTRIBUTES (decl)
-      = tree_cons (get_identifier ("omp target entrypoint"),
-                   NULL_TREE, DECL_ATTRIBUTES (decl));
+    {
+      const char *target_attr = (is_gimple_omp_offloaded (ctx->stmt)
+				 ? "omp target entrypoint"
+				 : "omp declare target");
+      DECL_ATTRIBUTES (decl)
+	= tree_cons (get_identifier (target_attr),
+		     NULL_TREE, DECL_ATTRIBUTES (decl));
+    }
 
   t = build_decl (DECL_SOURCE_LOCATION (decl),
 		  RESULT_DECL, NULL_TREE, void_type_node);
@@ -4264,6 +4274,25 @@ omp_clause_aligned_alignment (tree clause)
   return build_int_cst (integer_type_node, al);
 }
 
+
+/* Return maximum SIMT width if offloading may target SIMT hardware.  */
+
+static int
+omp_max_simt_vf (void)
+{
+  if (!optimize)
+    return 0;
+  if (ENABLE_OFFLOADING)
+    for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c; )
+      {
+	if (!strncmp (c, "nvptx", strlen ("nvptx")))
+	  return 32;
+	else if ((c = strchr (c, ',')))
+	  c++;
+      }
+  return 0;
+}
+
 /* Return maximum possible vectorization factor for the target.  */
 
 static int
@@ -4277,16 +4306,18 @@ omp_max_vf (void)
               || global_options_set.x_flag_tree_vectorize)))
     return 1;
 
+  int vf = 1;
   int vs = targetm.vectorize.autovectorize_vector_sizes ();
   if (vs)
+    vf = 1 << floor_log2 (vs);
+  else
     {
-      vs = 1 << floor_log2 (vs);
-      return vs;
+      machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
+      if (GET_MODE_CLASS (vqimode) == MODE_VECTOR_INT)
+	vf = GET_MODE_NUNITS (vqimode);
     }
-  machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
-  if (GET_MODE_CLASS (vqimode) == MODE_VECTOR_INT)
-    return GET_MODE_NUNITS (vqimode);
-  return 1;
+  int svf = omp_max_simt_vf ();
+  return MAX (vf, svf);
 }
 
 /* Helper function of lower_rec_input_clauses, used for #pragma omp simd
@@ -4374,10 +4405,13 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
   int pass;
   bool is_simd = (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
 		  && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD);
+  bool maybe_simt
+    = is_simd && omp_maybe_offloaded_ctx (ctx) && omp_max_simt_vf () > 1;
   int max_vf = 0;
   tree lane = NULL_TREE, idx = NULL_TREE;
+  tree simt_lane = NULL_TREE;
   tree ivar = NULL_TREE, lvar = NULL_TREE;
-  gimple_seq llist[2] = { NULL, NULL };
+  gimple_seq llist[3] = { };
 
   copyin_seq = NULL;
 
@@ -5251,6 +5285,16 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
 
 		      gimplify_assign (unshare_expr (ivar), x, &llist[0]);
 
+		      if (maybe_simt)
+			{
+			  if (!simt_lane)
+			    simt_lane = create_tmp_var (unsigned_type_node);
+			  x = build_call_expr_internal_loc
+			    (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_BFLY,
+			     TREE_TYPE (ivar), 2, ivar, simt_lane);
+			  x = build2 (code, TREE_TYPE (ivar), ivar, x);
+			  gimplify_assign (ivar, x, &llist[2]);
+			}
 		      x = build2 (code, TREE_TYPE (ref), ref, ivar);
 		      ref = build_outer_var_ref (var, ctx);
 		      gimplify_assign (ref, x, &llist[1]);
@@ -5303,6 +5347,39 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
       g = gimple_build_assign (lane, INTEGER_CST,
 			       build_int_cst (unsigned_type_node, 0));
       gimple_seq_add_stmt (ilist, g);
+      /* Emit reductions across SIMT lanes in log_2(simt_vf) steps.  */
+      if (llist[2])
+	{
+	  tree simt_vf = create_tmp_var (unsigned_type_node);
+	  g = gimple_build_call_internal (IFN_GOMP_SIMT_VF, 0);
+	  gimple_call_set_lhs (g, simt_vf);
+	  gimple_seq_add_stmt (dlist, g);
+
+	  tree t = build_int_cst (unsigned_type_node, 1);
+	  g = gimple_build_assign (simt_lane, INTEGER_CST, t);
+	  gimple_seq_add_stmt (dlist, g);
+
+	  t = build_int_cst (unsigned_type_node, 0);
+	  g = gimple_build_assign (idx, INTEGER_CST, t);
+	  gimple_seq_add_stmt (dlist, g);
+
+	  tree body = create_artificial_label (UNKNOWN_LOCATION);
+	  tree header = create_artificial_label (UNKNOWN_LOCATION);
+	  tree end = create_artificial_label (UNKNOWN_LOCATION);
+	  gimple_seq_add_stmt (dlist, gimple_build_goto (header));
+	  gimple_seq_add_stmt (dlist, gimple_build_label (body));
+
+	  gimple_seq_add_seq (dlist, llist[2]);
+
+	  g = gimple_build_assign (simt_lane, LSHIFT_EXPR, simt_lane, integer_one_node);
+	  gimple_seq_add_stmt (dlist, g);
+
+	  gimple_seq_add_stmt (dlist, gimple_build_label (header));
+	  g = gimple_build_cond (LT_EXPR, simt_lane, simt_vf, body, end);
+	  gimple_seq_add_stmt (dlist, g);
+
+	  gimple_seq_add_stmt (dlist, gimple_build_label (end));
+	}
       for (int i = 0; i < 2; i++)
 	if (llist[i])
 	  {
@@ -5389,7 +5466,7 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
 {
   tree x, c, label = NULL, orig_clauses = clauses;
   bool par_clauses = false;
-  tree simduid = NULL, lastlane = NULL;
+  tree simduid = NULL, lastlane = NULL, simtcond = NULL, simtlast = NULL;
 
   /* Early exit if there are no lastprivate or linear clauses.  */
   for (; clauses ; clauses = OMP_CLAUSE_CHAIN (clauses))
@@ -5416,6 +5493,16 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
       par_clauses = true;
     }
 
+  bool maybe_simt = false;
+  if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
+      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
+    {
+      maybe_simt = omp_maybe_offloaded_ctx (ctx) && omp_max_simt_vf () > 1;
+      simduid = find_omp_clause (orig_clauses, OMP_CLAUSE__SIMDUID_);
+      if (simduid)
+	simduid = OMP_CLAUSE__SIMDUID__DECL (simduid);
+    }
+
   if (predicate)
     {
       gcond *stmt;
@@ -5427,20 +5514,27 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
       arm2 = TREE_OPERAND (predicate, 1);
       gimplify_expr (&arm1, stmt_list, NULL, is_gimple_val, fb_rvalue);
       gimplify_expr (&arm2, stmt_list, NULL, is_gimple_val, fb_rvalue);
-      stmt = gimple_build_cond (TREE_CODE (predicate), arm1, arm2,
-				label_true, label);
+      if (maybe_simt)
+	{
+	  c = build2 (TREE_CODE (predicate), boolean_type_node, arm1, arm2);
+	  c = fold_convert (integer_type_node, c);
+	  simtcond = create_tmp_var (integer_type_node);
+	  gimplify_assign (simtcond, c, stmt_list);
+	  gcall *g = gimple_build_call_internal (IFN_GOMP_SIMT_VOTE_ANY,
+						 1, simtcond);
+	  c = create_tmp_var (integer_type_node);
+	  gimple_call_set_lhs (g, c);
+	  gimple_seq_add_stmt (stmt_list, g);
+	  stmt = gimple_build_cond (NE_EXPR, c, integer_zero_node,
+				    label_true, label);
+	}
+      else
+	stmt = gimple_build_cond (TREE_CODE (predicate), arm1, arm2,
+				  label_true, label);
       gimple_seq_add_stmt (stmt_list, stmt);
       gimple_seq_add_stmt (stmt_list, gimple_build_label (label_true));
     }
 
-  if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
-      && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
-    {
-      simduid = find_omp_clause (orig_clauses, OMP_CLAUSE__SIMDUID_);
-      if (simduid)
-	simduid = OMP_CLAUSE__SIMDUID__DECL (simduid);
-    }
-
   for (c = clauses; c ;)
     {
       tree var, new_var;
@@ -5491,6 +5585,24 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
 		  new_var = build4 (ARRAY_REF, TREE_TYPE (val),
 				    TREE_OPERAND (val, 0), lastlane,
 				    NULL_TREE, NULL_TREE);
+		  if (maybe_simt)
+		    {
+		      gcall *g;
+		      if (simtlast == NULL)
+			{
+			  simtlast = create_tmp_var (unsigned_type_node);
+			  g = gimple_build_call_internal
+			    (IFN_GOMP_SIMT_LAST_LANE, 1, simtcond);
+			  gimple_call_set_lhs (g, simtlast);
+			  gimple_seq_add_stmt (stmt_list, g);
+			}
+		      x = build_call_expr_internal_loc
+			(UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX,
+			 TREE_TYPE (new_var), 2, new_var, simtlast);
+		      new_var = unshare_expr (new_var);
+		      gimplify_assign (new_var, x, stmt_list);
+		      new_var = unshare_expr (new_var);
+		    }
 		}
 	    }
 
@@ -10564,12 +10676,23 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd)
   edge e, ne;
   tree *counts = NULL;
   int i;
+  int safelen_int = INT_MAX;
   tree safelen = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt),
 				  OMP_CLAUSE_SAFELEN);
   tree simduid = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt),
 				  OMP_CLAUSE__SIMDUID_);
   tree n1, n2;
 
+  if (safelen)
+    {
+      safelen = OMP_CLAUSE_SAFELEN_EXPR (safelen);
+      if (TREE_CODE (safelen) != INTEGER_CST)
+	safelen_int = 0;
+      else if (tree_fits_uhwi_p (safelen) && tree_to_uhwi (safelen) < INT_MAX)
+	safelen_int = tree_to_uhwi (safelen);
+      if (safelen_int == 1)
+	safelen_int = 0;
+    }
   type = TREE_TYPE (fd->loop.v);
   entry_bb = region->entry;
   cont_bb = region->cont;
@@ -10623,20 +10746,53 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd)
 				OMP_CLAUSE__LOOPTEMP_);
       gcc_assert (innerc);
       n2 = OMP_CLAUSE_DECL (innerc);
-      expand_omp_build_assign (&gsi, fd->loop.v,
-			       fold_convert (type, n1));
+    }
+  tree step = fd->loop.step;
+
+  bool offloaded = cgraph_node::get (current_function_decl)->offloadable;
+  for (struct omp_region *rgn = region; !offloaded && rgn; rgn = rgn->outer)
+    offloaded = rgn->type == GIMPLE_OMP_TARGET;
+  bool is_simt = offloaded && omp_max_simt_vf () > 1 && safelen_int > 1;
+  tree simt_lane = NULL_TREE, simt_maxlane = NULL_TREE;
+  if (is_simt)
+    {
+      cfun->curr_properties &= ~PROP_gimple_lomp_dev;
+      simt_lane = create_tmp_var (unsigned_type_node);
+      gimple *g = gimple_build_call_internal (IFN_GOMP_SIMT_LANE, 0);
+      gimple_call_set_lhs (g, simt_lane);
+      gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+      tree offset = fold_build2 (MULT_EXPR, TREE_TYPE (step), step,
+				 fold_convert (TREE_TYPE (step), simt_lane));
+      n1 = fold_convert (type, n1);
+      if (POINTER_TYPE_P (type))
+	n1 = fold_build_pointer_plus (n1, offset);
+      else
+	n1 = fold_build2 (PLUS_EXPR, type, n1, fold_convert (type, offset));
+
+      /* Collapsed loops not handled for SIMT yet: limit to one lane only.  */
       if (fd->collapse > 1)
+	simt_maxlane = build_one_cst (unsigned_type_node);
+      else if (safelen_int < omp_max_simt_vf ())
+	simt_maxlane = build_int_cst (unsigned_type_node, safelen_int);
+      tree vf
+	= build_call_expr_internal_loc (UNKNOWN_LOCATION, IFN_GOMP_SIMT_VF,
+					unsigned_type_node, 0);
+      if (simt_maxlane)
+	vf = fold_build2 (MIN_EXPR, unsigned_type_node, vf, simt_maxlane);
+      vf = fold_convert (TREE_TYPE (step), vf);
+      step = fold_build2 (MULT_EXPR, TREE_TYPE (step), step, vf);
+    }
+
+  expand_omp_build_assign (&gsi, fd->loop.v, fold_convert (type, n1));
+  if (fd->collapse > 1)
+    {
+      if (gimple_omp_for_combined_into_p (fd->for_stmt))
 	{
 	  gsi_prev (&gsi);
 	  expand_omp_for_init_vars (fd, &gsi, counts, NULL, n1);
 	  gsi_next (&gsi);
 	}
-    }
-  else
-    {
-      expand_omp_build_assign (&gsi, fd->loop.v,
-			       fold_convert (type, fd->loop.n1));
-      if (fd->collapse > 1)
+      else
 	for (i = 0; i < fd->collapse; i++)
 	  {
 	    tree itype = TREE_TYPE (fd->loops[i].v);
@@ -10645,7 +10801,7 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd)
 	    t = fold_convert (TREE_TYPE (fd->loops[i].v), fd->loops[i].n1);
 	    expand_omp_build_assign (&gsi, fd->loops[i].v, t);
 	  }
-      }
+    }
 
   /* Remove the GIMPLE_OMP_FOR statement.  */
   gsi_remove (&gsi, true);
@@ -10658,9 +10814,9 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd)
       gcc_assert (gimple_code (stmt) == GIMPLE_OMP_CONTINUE);
 
       if (POINTER_TYPE_P (type))
-	t = fold_build_pointer_plus (fd->loop.v, fd->loop.step);
+	t = fold_build_pointer_plus (fd->loop.v, step);
       else
-	t = fold_build2 (PLUS_EXPR, type, fd->loop.v, fd->loop.step);
+	t = fold_build2 (PLUS_EXPR, type, fd->loop.v, step);
       expand_omp_build_assign (&gsi, fd->loop.v, t);
 
       if (fd->collapse > 1)
@@ -10734,6 +10890,18 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd)
       gimple_regimplify_operands (cond_stmt, &gsi);
     }
 
+  /* Add 'V -= STEP * (SIMT_VF - 1)' after the loop.  */
+  if (is_simt)
+    {
+      gsi = gsi_start_bb (l2_bb);
+      step = fold_build2 (MINUS_EXPR, TREE_TYPE (step), fd->loop.step, step);
+      if (POINTER_TYPE_P (type))
+	t = fold_build_pointer_plus (fd->loop.v, step);
+      else
+	t = fold_build2 (PLUS_EXPR, type, fd->loop.v, step);
+      expand_omp_build_assign (&gsi, fd->loop.v, t);
+    }
+
   /* Remove GIMPLE_OMP_RETURN.  */
   gsi = gsi_last_bb (exit_bb);
   gsi_remove (&gsi, true);
@@ -10763,30 +10931,29 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd)
   ne->probability = REG_BR_PROB_BASE / 8;
 
   set_immediate_dominator (CDI_DOMINATORS, l1_bb, entry_bb);
-  set_immediate_dominator (CDI_DOMINATORS, l2_bb, l2_dom_bb);
   set_immediate_dominator (CDI_DOMINATORS, l0_bb, l1_bb);
 
+  if (simt_maxlane)
+    {
+      cond_stmt = gimple_build_cond (LT_EXPR, simt_lane, simt_maxlane,
+				     NULL_TREE, NULL_TREE);
+      gsi = gsi_last_bb (entry_bb);
+      gsi_insert_after (&gsi, cond_stmt, GSI_NEW_STMT);
+      make_edge (entry_bb, l2_bb, EDGE_FALSE_VALUE);
+      FALLTHRU_EDGE (entry_bb)->flags = EDGE_TRUE_VALUE;
+      FALLTHRU_EDGE (entry_bb)->probability = REG_BR_PROB_BASE * 7 / 8;
+      BRANCH_EDGE (entry_bb)->probability = REG_BR_PROB_BASE / 8;
+      l2_dom_bb = entry_bb;
+    }
+  set_immediate_dominator (CDI_DOMINATORS, l2_bb, l2_dom_bb);
+
   if (!broken_loop)
     {
       struct loop *loop = alloc_loop ();
       loop->header = l1_bb;
       loop->latch = cont_bb;
       add_loop (loop, l1_bb->loop_father);
-      if (safelen == NULL_TREE)
-	loop->safelen = INT_MAX;
-      else
-	{
-	  safelen = OMP_CLAUSE_SAFELEN_EXPR (safelen);
-	  if (TREE_CODE (safelen) != INTEGER_CST)
-	    loop->safelen = 0;
-	  else if (!tree_fits_uhwi_p (safelen)
-		   || tree_to_uhwi (safelen) > INT_MAX)
-	    loop->safelen = INT_MAX;
-	  else
-	    loop->safelen = tree_to_uhwi (safelen);
-	  if (loop->safelen == 1)
-	    loop->safelen = 0;
-	}
+      loop->safelen = safelen_int;
       if (simduid)
 	{
 	  loop->simduid = OMP_CLAUSE__SIMDUID__DECL (simduid);
@@ -13951,7 +14118,6 @@ expand_omp (struct omp_region *region)
     }
 }
 
-
 /* Helper for build_omp_regions.  Scan the dominator tree starting at
    block BB.  PARENT is the region that contains BB.  If SINGLE_TREE is
    true, the function ends once a single tree is built (otherwise, whole
@@ -14834,12 +15000,14 @@ static void
 lower_omp_ordered (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 {
   tree block;
-  gimple *stmt = gsi_stmt (*gsi_p);
+  gimple *stmt = gsi_stmt (*gsi_p), *g;
   gomp_ordered *ord_stmt = as_a <gomp_ordered *> (stmt);
   gcall *x;
   gbind *bind;
   bool simd = find_omp_clause (gimple_omp_ordered_clauses (ord_stmt),
 			       OMP_CLAUSE_SIMD);
+  bool maybe_simt
+    = simd && omp_maybe_offloaded_ctx (ctx) && omp_max_simt_vf () > 1;
   bool threads = find_omp_clause (gimple_omp_ordered_clauses (ord_stmt),
 				  OMP_CLAUSE_THREADS);
 
@@ -14873,11 +15041,56 @@ lower_omp_ordered (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 			   0);
   gimple_bind_add_stmt (bind, x);
 
+  tree counter = NULL_TREE, test = NULL_TREE, body = NULL_TREE;
+  if (maybe_simt)
+    {
+      counter = create_tmp_var (integer_type_node);
+      g = gimple_build_call_internal (IFN_GOMP_SIMT_LANE, 0);
+      gimple_call_set_lhs (g, counter);
+      gimple_bind_add_stmt (bind, g);
+
+      body = create_artificial_label (UNKNOWN_LOCATION);
+      test = create_artificial_label (UNKNOWN_LOCATION);
+      gimple_bind_add_stmt (bind, gimple_build_label (body));
+
+      tree simt_pred = create_tmp_var (integer_type_node);
+      g = gimple_build_call_internal (IFN_GOMP_SIMT_ORDERED_PRED, 1, counter);
+      gimple_call_set_lhs (g, simt_pred);
+      gimple_bind_add_stmt (bind, g);
+
+      tree t = create_artificial_label (UNKNOWN_LOCATION);
+      g = gimple_build_cond (EQ_EXPR, simt_pred, integer_zero_node, t, test);
+      gimple_bind_add_stmt (bind, g);
+
+      gimple_bind_add_stmt (bind, gimple_build_label (t));
+    }
   lower_omp (gimple_omp_body_ptr (stmt), ctx);
   gimple_omp_set_body (stmt, maybe_catch_exception (gimple_omp_body (stmt)));
   gimple_bind_add_seq (bind, gimple_omp_body (stmt));
   gimple_omp_set_body (stmt, NULL);
 
+  if (maybe_simt)
+    {
+      gimple_bind_add_stmt (bind, gimple_build_label (test));
+      g = gimple_build_assign (counter, MINUS_EXPR, counter, integer_one_node);
+      gimple_bind_add_stmt (bind, g);
+
+      tree c = build2 (GE_EXPR, boolean_type_node, counter, integer_zero_node);
+      tree nonneg = create_tmp_var (integer_type_node);
+      gimple_seq tseq = NULL;
+      gimplify_assign (nonneg, fold_convert (integer_type_node, c), &tseq);
+      gimple_bind_add_seq (bind, tseq);
+
+      g = gimple_build_call_internal (IFN_GOMP_SIMT_VOTE_ANY, 1, nonneg);
+      gimple_call_set_lhs (g, nonneg);
+      gimple_bind_add_stmt (bind, g);
+
+      tree end = create_artificial_label (UNKNOWN_LOCATION);
+      g = gimple_build_cond (NE_EXPR, nonneg, integer_zero_node, body, end);
+      gimple_bind_add_stmt (bind, g);
+
+      gimple_bind_add_stmt (bind, gimple_build_label (end));
+    }
   if (simd)
     x = gimple_build_call_internal (IFN_GOMP_SIMD_ORDERED_END, 1,
 				    build_int_cst (NULL_TREE, threads));
@@ -17998,7 +18211,7 @@ const pass_data pass_data_lower_omp =
   OPTGROUP_NONE, /* optinfo_flags */
   TV_NONE, /* tv_id */
   PROP_gimple_any, /* properties_required */
-  PROP_gimple_lomp, /* properties_provided */
+  PROP_gimple_lomp | PROP_gimple_lomp_dev, /* properties_provided */
   0, /* properties_destroyed */
   0, /* todo_flags_start */
   0, /* todo_flags_finish */
@@ -19930,6 +20143,113 @@ make_pass_oacc_device_lower (gcc::context *ctxt)
 {
   return new pass_oacc_device_lower (ctxt);
 }
+
+
+/* Cleanup uses of SIMT placeholder internal functions: on non-SIMT targets,
+   VF is 1 and LANE is 0; on SIMT targets, VF is folded to a constant, and
+   LANE is kept to be expanded to RTL later on.  Also cleanup all other SIMT
+   internal functions on non-SIMT targets, and likewise some SIMD internal
+   functions on SIMT targets.  */
+
+static unsigned int
+execute_omp_device_lower ()
+{
+  int vf = targetm.simt.vf ? targetm.simt.vf () : 1;
+  basic_block bb;
+  gimple_stmt_iterator gsi;
+  FOR_EACH_BB_FN (bb, cfun)
+    for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
+      {
+	gimple *stmt = gsi_stmt (gsi);
+	if (!is_gimple_call (stmt) || !gimple_call_internal_p (stmt))
+	  continue;
+	tree lhs = gimple_call_lhs (stmt), rhs = NULL_TREE;
+	tree type = lhs ? TREE_TYPE (lhs) : integer_type_node;
+	switch (gimple_call_internal_fn (stmt))
+	  {
+	  case IFN_GOMP_SIMT_LANE:
+	  case IFN_GOMP_SIMT_LAST_LANE:
+	    rhs = vf == 1 ? build_zero_cst (type) : NULL_TREE;
+	    break;
+	  case IFN_GOMP_SIMT_VF:
+	    rhs = build_int_cst (type, vf);
+	    break;
+	  case IFN_GOMP_SIMT_ORDERED_PRED:
+	    rhs = vf == 1 ? integer_zero_node : NULL_TREE;
+	    if (rhs || !lhs)
+	      unlink_stmt_vdef (stmt);
+	    break;
+	  case IFN_GOMP_SIMT_VOTE_ANY:
+	  case IFN_GOMP_SIMT_XCHG_BFLY:
+	  case IFN_GOMP_SIMT_XCHG_IDX:
+	    rhs = vf == 1 ? gimple_call_arg (stmt, 0) : NULL_TREE;
+	    break;
+	  case IFN_GOMP_SIMD_LANE:
+	  case IFN_GOMP_SIMD_LAST_LANE:
+	    rhs = vf != 1 ? build_zero_cst (type) : NULL_TREE;
+	    break;
+	  case IFN_GOMP_SIMD_VF:
+	    rhs = vf != 1 ? build_one_cst (type) : NULL_TREE;
+	    break;
+	  default:
+	    continue;
+	  }
+	if (lhs && !rhs)
+	  continue;
+	stmt = lhs ? gimple_build_assign (lhs, rhs) : gimple_build_nop ();
+	gsi_replace (&gsi, stmt, false);
+      }
+  if (vf != 1)
+    cfun->has_force_vectorize_loops = false;
+  return 0;
+}
+
+namespace {
+
+const pass_data pass_data_omp_device_lower =
+{
+  GIMPLE_PASS, /* type */
+  "ompdevlow", /* name */
+  OPTGROUP_NONE, /* optinfo_flags */
+  TV_NONE, /* tv_id */
+  PROP_cfg, /* properties_required */
+  PROP_gimple_lomp_dev, /* properties_provided */
+  0, /* properties_destroyed */
+  0, /* todo_flags_start */
+  TODO_update_ssa, /* todo_flags_finish */
+};
+
+class pass_omp_device_lower : public gimple_opt_pass
+{
+public:
+  pass_omp_device_lower (gcc::context *ctxt)
+    : gimple_opt_pass (pass_data_omp_device_lower, ctxt)
+  {}
+
+  /* opt_pass methods: */
+  virtual bool gate (function *ARG_UNUSED (fun))
+    {
+      /* FIXME: this should use PROP_gimple_lomp_dev.  */
+#ifdef ACCEL_COMPILER
+      return true;
+#else
+      return ENABLE_OFFLOADING && (flag_openmp || in_lto_p);
+#endif
+    }
+  virtual unsigned int execute (function *)
+    {
+      return execute_omp_device_lower ();
+    }
+
+}; // class pass_expand_omp_ssa
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_omp_device_lower (gcc::context *ctxt)
+{
+  return new pass_omp_device_lower (ctxt);
+}
 
 /* "omp declare target link" handling pass.  */