Decompose OpenACC 'kernels' constructs into parts, a sequence of compute constructs

Not yet enabled by default: for now, the current mode of OpenACC 'kernels'
constructs handling still remains '-fopenacc-kernels=parloops', but that is to
change later.

	gcc/
	* omp-oacc-kernels-decompose.cc: New.
	* Makefile.in (OBJS): Add it.
	* passes.def: Instantiate it.
	* tree-pass.h (make_pass_omp_oacc_kernels_decompose): Declare.
	* flag-types.h (enum openacc_kernels): Add.
	* doc/invoke.texi (-fopenacc-kernels): Document.
	* gimple.h (enum gf_mask): Add
	'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED',
	'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE',
	'GF_OMP_TARGET_KIND_OACC_DATA_KERNELS'.
	(is_gimple_omp_oacc, is_gimple_omp_offloaded): Handle these.
	* gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
	* omp-expand.c (expand_omp_target, build_omp_regions_1)
	(omp_make_gimple_edges): Likewise.
	* omp-low.c (scan_sharing_clauses, scan_omp_for)
	(check_omp_nesting_restrictions, lower_oacc_reductions)
	(lower_oacc_head_mark, lower_omp_target): Likewise.
	* omp-offload.c (execute_oacc_device_lower): Likewise.
	gcc/c-family/
	* c.opt (fopenacc-kernels): Add.
	gcc/fortran/
	* lang.opt (fopenacc-kernels): Add.
	gcc/testsuite/
	* c-c++-common/goacc/kernels-decompose-1.c: New.
	* c-c++-common/goacc/kernels-decompose-2.c: New.
	* c-c++-common/goacc/kernels-decompose-ice-1.c: New.
	* c-c++-common/goacc/kernels-decompose-ice-2.c: New.
	* gfortran.dg/goacc/kernels-decompose-1.f95: New.
	* gfortran.dg/goacc/kernels-decompose-2.f95: New.
	* c-c++-common/goacc/if-clause-2.c: Adjust.
	* gfortran.dg/goacc/kernels-tree.f95: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c:
	New.
	* testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Adjust.
	* testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise.

Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>

1 files changed, 1531 insertions, 0 deletions

diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
new file mode 100644
index 0000000..c585e5d
--- /dev/null
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -0,0 +1,1531 @@
+/* Decompose OpenACC 'kernels' constructs into parts, a sequence of compute
+   constructs
+
+   Copyright (C) 2020 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "config.h"
+#include "system.h"
+#include "coretypes.h"
+#include "backend.h"
+#include "target.h"
+#include "tree.h"
+#include "cp/cp-tree.h"
+#include "gimple.h"
+#include "tree-pass.h"
+#include "cgraph.h"
+#include "fold-const.h"
+#include "gimplify.h"
+#include "gimple-iterator.h"
+#include "gimple-walk.h"
+#include "gomp-constants.h"
+#include "omp-general.h"
+#include "diagnostic-core.h"
+
+
+/* This preprocessing pass is run immediately before lower_omp.  It decomposes
+   OpenACC 'kernels' constructs into parts, a sequence of compute constructs.
+
+   The translation is as follows:
+     - The entire 'kernels' region is turned into a 'data' region with clauses
+       taken from the 'kernels' region.  New 'create' clauses are added for all
+       variables declared at the top level in the kernels region.
+     - Any loop nests annotated with an OpenACC 'loop' directive are wrapped in
+       a new compute construct.
+	 - 'loop' directives without an explicit 'independent' or 'seq' clause
+	   get an 'auto' clause added; other clauses are preserved on the loop
+	   or moved to the new surrounding compute construct, as applicable.
+     - Any sequences of other code (non-loops, non-OpenACC 'loop's) are wrapped
+       in new "gang-single" compute construct: 'worker'/'vector' parallelism is
+       preserved, but 'num_gangs (1)' is enforced.
+     - Both points above only apply at the topmost level in the region, that
+       is, the transformation does not introduce new compute constructs inside
+       nested statement bodies.  In particular, this means that a
+       gang-parallelizable loop inside an 'if' statement is made "gang-single".
+     - In order to make the host wait only once for the whole region instead
+       of once per device kernel launch, the new compute constructs are
+       annotated 'async'.  Unless the original 'kernels' construct already was
+       marked 'async', the entire region ends with a 'wait' directive.  If the
+       original 'kernels' construct was marked 'async', the synthesized 'async'
+       clauses use the original 'kernels' construct's 'async' argument
+       (possibly implicit).
+*/
+
+
+/*TODO Things are conceptually wrong here: 'loop' clauses may be hidden behind
+  'device_type', so we have to defer a lot of processing until we're in the
+  offloading compilation.  "Fortunately", GCC doesn't support the OpenACC
+  'device_type' clause yet, so we get away that.  */
+
+
+/* Helper function for decompose_kernels_region_body.  If STMT contains a
+   "top-level" OMP_FOR statement, returns a pointer to that statement;
+   returns NULL otherwise.
+
+   A "top-level" OMP_FOR statement is one that is possibly accompanied by
+   small snippets of setup code.  Specifically, this function accepts an
+   OMP_FOR possibly wrapped in a singleton bind and a singleton try
+   statement to allow for a local loop variable, but not an OMP_FOR
+   statement nested in any other constructs.  Alternatively, it accepts a
+   non-singleton bind containing only assignments and then an OMP_FOR
+   statement at the very end.  The former style can be generated by the C
+   frontend, the latter by the Fortran frontend.  */
+
+static gimple *
+top_level_omp_for_in_stmt (gimple *stmt)
+{
+  if (gimple_code (stmt) == GIMPLE_OMP_FOR)
+    return stmt;
+
+  if (gimple_code (stmt) == GIMPLE_BIND)
+    {
+      gimple_seq body = gimple_bind_body (as_a <gbind *> (stmt));
+      if (gimple_seq_singleton_p (body))
+	{
+	  /* Accept an OMP_FOR statement, or a try statement containing only
+	     a single OMP_FOR.  */
+	  gimple *maybe_for_or_try = gimple_seq_first_stmt (body);
+	  if (gimple_code (maybe_for_or_try) == GIMPLE_OMP_FOR)
+	    return maybe_for_or_try;
+	  else if (gimple_code (maybe_for_or_try) == GIMPLE_TRY)
+	    {
+	      gimple_seq try_body = gimple_try_eval (maybe_for_or_try);
+	      if (!gimple_seq_singleton_p (try_body))
+		return NULL;
+	      gimple *maybe_omp_for_stmt = gimple_seq_first_stmt (try_body);
+	      if (gimple_code (maybe_omp_for_stmt) == GIMPLE_OMP_FOR)
+		return maybe_omp_for_stmt;
+	    }
+	}
+      else
+	{
+	  gimple_stmt_iterator gsi;
+	  /* Accept only a block of optional assignments followed by an
+	     OMP_FOR at the end.  No other kinds of statements allowed.  */
+	  for (gsi = gsi_start (body); !gsi_end_p (gsi); gsi_next (&gsi))
+	    {
+	      gimple *body_stmt = gsi_stmt (gsi);
+	      if (gimple_code (body_stmt) == GIMPLE_ASSIGN)
+		continue;
+	      else if (gimple_code (body_stmt) == GIMPLE_OMP_FOR
+		       && gsi_one_before_end_p (gsi))
+		return body_stmt;
+	      else
+		return NULL;
+	    }
+	}
+    }
+
+  return NULL;
+}
+
+/* Helper for adjust_region_code: evaluate the statement at GSI_P.  */
+
+static tree
+adjust_region_code_walk_stmt_fn (gimple_stmt_iterator *gsi_p,
+				 bool *handled_ops_p,
+				 struct walk_stmt_info *wi)
+{
+  int *region_code = (int *) wi->info;
+
+  gimple *stmt = gsi_stmt (*gsi_p);
+  switch (gimple_code (stmt))
+    {
+    case GIMPLE_OMP_FOR:
+      {
+	tree clauses = gimple_omp_for_clauses (stmt);
+	if (omp_find_clause (clauses, OMP_CLAUSE_INDEPENDENT))
+	  {
+	    /* Explicit 'independent' clause.  */
+	    /* Keep going; recurse into loop body.  */
+	    break;
+	  }
+	else if (omp_find_clause (clauses, OMP_CLAUSE_SEQ))
+	  {
+	    /* Explicit 'seq' clause.  */
+	    /* We'll "parallelize" if at some level a loop construct has been
+	       marked up by the user as unparallelizable ('seq' clause; we'll
+	       respect that in the later processing).  Given that the user has
+	       explicitly marked it up, this loop construct cannot be
+	       performance-critical, and in this case it's also fine to
+	       "parallelize" instead of "gang-single", because any outer or
+	       inner loops may still exploit the available parallelism.  */
+	    /* Keep going; recurse into loop body.  */
+	    break;
+	  }
+	else
+	  {
+	    /* Explicit or implicit 'auto' clause.  */
+	    /* The user would like this loop analyzed ('auto' clause) and
+	       typically parallelized, but we don't have available yet the
+	       compiler logic to analyze this, so can't parallelize it here, so
+	       we'd very likely be running into a performance problem if we
+	       were to execute this unparallelized, thus forward the whole loop
+	       nest to 'parloops'.  */
+	    *region_code = GF_OMP_TARGET_KIND_OACC_KERNELS;
+	    /* Terminate: final decision for this region.  */
+	    *handled_ops_p = true;
+	    return integer_zero_node;
+	  }
+	gcc_unreachable ();
+      }
+
+    case GIMPLE_COND:
+    case GIMPLE_GOTO:
+    case GIMPLE_SWITCH:
+    case GIMPLE_ASM:
+    case GIMPLE_TRANSACTION:
+    case GIMPLE_RETURN:
+      /* Statement that might constitute some looping/control flow pattern.  */
+      /* The user would like this code analyzed (implicit inside a 'kernels'
+	 region) and typically parallelized, but we don't have available yet
+	 the compiler logic to analyze this, so can't parallelize it here, so
+	 we'd very likely be running into a performance problem if we were to
+	 execute this unparallelized, thus forward the whole thing to
+	 'parloops'.  */
+      *region_code = GF_OMP_TARGET_KIND_OACC_KERNELS;
+      /* Terminate: final decision for this region.  */
+      *handled_ops_p = true;
+      return integer_zero_node;
+
+    default:
+      /* Keep going.  */
+      break;
+    }
+
+  return NULL;
+}
+
+/* Adjust the REGION_CODE for the region in GS.  */
+
+static void
+adjust_region_code (gimple_seq gs, int *region_code)
+{
+  struct walk_stmt_info wi;
+  memset (&wi, 0, sizeof (wi));
+  wi.info = region_code;
+  walk_gimple_seq (gs, adjust_region_code_walk_stmt_fn, NULL, &wi);
+}
+
+/* Helper function for make_loops_gang_single for walking the tree.  If the
+   statement indicated by GSI_P is an OpenACC for loop with a gang clause,
+   issue a warning and remove the clause.  */
+
+static tree
+visit_loops_in_gang_single_region (gimple_stmt_iterator *gsi_p,
+				   bool *handled_ops_p,
+				   struct walk_stmt_info *)
+{
+  *handled_ops_p = false;
+
+  gimple *stmt = gsi_stmt (*gsi_p);
+  switch (gimple_code (stmt))
+    {
+    case GIMPLE_OMP_FOR:
+      /*TODO Given the current 'adjust_region_code' algorithm, this is
+	actually...  */
+      gcc_unreachable ();
+
+      {
+	tree clauses = gimple_omp_for_clauses (stmt);
+	tree prev_clause = NULL;
+	for (tree clause = clauses; clause; clause = OMP_CLAUSE_CHAIN (clause))
+	  {
+	    if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_GANG)
+	      {
+		/* It makes no sense to have a 'gang' clause in a "gang-single"
+		   region, so warn and remove it.  */
+		warning_at (gimple_location (stmt), 0,
+			    "conditionally executed loop in %<kernels%> region"
+			    " will be executed by a single gang;"
+			    " ignoring %<gang%> clause");
+		if (prev_clause != NULL)
+		  OMP_CLAUSE_CHAIN (prev_clause) = OMP_CLAUSE_CHAIN (clause);
+		else
+		  clauses = OMP_CLAUSE_CHAIN (clause);
+
+		break;
+	      }
+	    prev_clause = clause;
+	  }
+	gimple_omp_for_set_clauses (stmt, clauses);
+      }
+      /* No need to recurse into nested statements; no loop nested inside
+	 this loop can be gang-partitioned.  */
+      sorry ("%<gang%> loop in %<gang-single%> region");
+      *handled_ops_p = true;
+      break;
+
+    default:
+      break;
+    }
+
+  return NULL;
+}
+
+/* Visit all nested OpenACC loops in the sequence indicated by GS.  This
+   statement is expected to be inside a gang-single region.  Issue a warning
+   for any loops inside it that have gang clauses and remove the clauses.  */
+
+static void
+make_loops_gang_single (gimple_seq gs)
+{
+  struct walk_stmt_info wi;
+  memset (&wi, 0, sizeof (wi));
+  walk_gimple_seq (gs, visit_loops_in_gang_single_region, NULL, &wi);
+}
+
+/* Construct a "gang-single" compute construct at LOC containing the STMTS.
+   Annotate with CLAUSES, which must not contain a 'num_gangs' clause, and an
+   additional 'num_gangs (1)' clause to force "gang-single" execution.  */
+
+static gimple *
+make_region_seq (location_t loc, gimple_seq stmts,
+		 tree num_gangs_clause,
+		 tree num_workers_clause,
+		 tree vector_length_clause,
+		 tree clauses)
+{
+  /* This correctly unshares the entire clause chain rooted here.  */
+  clauses = unshare_expr (clauses);
+
+  dump_user_location_t loc_stmts_first = gimple_seq_first (stmts);
+
+  /* Figure out the region code for this region.  */
+  /* Optimistic default: assume "setup code", no looping; thus not
+     performance-critical.  */
+  int region_code = GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE;
+  adjust_region_code (stmts, &region_code);
+
+  if (region_code == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE)
+    {
+      if (dump_enabled_p ())
+	/*TODO MSG_MISSED_OPTIMIZATION? */
+	dump_printf_loc (MSG_NOTE, loc_stmts_first,
+			 "beginning %<gang-single%> part"
+			 " in OpenACC %<kernels%> region\n");
+
+      /* Synthesize a 'num_gangs (1)' clause.  */
+      tree gang_single_clause = build_omp_clause (loc, OMP_CLAUSE_NUM_GANGS);
+      OMP_CLAUSE_OPERAND (gang_single_clause, 0) = integer_one_node;
+      OMP_CLAUSE_CHAIN (gang_single_clause) = clauses;
+      clauses = gang_single_clause;
+
+      /* Remove and issue warnings about gang clauses on any OpenACC
+	 loops nested inside this sequentially executed statement.  */
+      make_loops_gang_single (stmts);
+    }
+  else if (region_code == GF_OMP_TARGET_KIND_OACC_KERNELS)
+    {
+      if (dump_enabled_p ())
+	dump_printf_loc (MSG_NOTE, loc_stmts_first,
+			 "beginning %<parloops%> part"
+			 " in OpenACC %<kernels%> region\n");
+
+      /* As we're transforming a 'GF_OMP_TARGET_KIND_OACC_KERNELS' into another
+	 'GF_OMP_TARGET_KIND_OACC_KERNELS', this isn't doing any of the clauses
+	 mangling that 'make_region_loop_nest' is doing.  */
+      /* Re-assemble the clauses stripped off earlier.  */
+      if (num_gangs_clause != NULL)
+	{
+	  tree c = unshare_expr (num_gangs_clause);
+	  OMP_CLAUSE_CHAIN (c) = clauses;
+	  clauses = c;
+	}
+      if (num_workers_clause != NULL)
+	{
+	  tree c = unshare_expr (num_workers_clause);
+	  OMP_CLAUSE_CHAIN (c) = clauses;
+	  clauses = c;
+	}
+      if (vector_length_clause != NULL)
+	{
+	  tree c = unshare_expr (vector_length_clause);
+	  OMP_CLAUSE_CHAIN (c) = clauses;
+	  clauses = c;
+	}
+    }
+  else
+    gcc_unreachable ();
+
+  /* Build the gang-single region.  */
+  gimple *single_region = gimple_build_omp_target (NULL, region_code, clauses);
+  gimple_set_location (single_region, loc);
+  gbind *single_body = gimple_build_bind (NULL, stmts, make_node (BLOCK));
+  gimple_omp_set_body (single_region, single_body);
+
+  return single_region;
+}
+
+/* Helper function for make_region_loop_nest.  Adds a 'num_gangs'
+   ('num_workers', 'vector_length') clause to the given CLAUSES, either the one
+   from the parent compute construct (PARENT_CLAUSE) or a new one based on the
+   loop's own LOOP_CLAUSE ('gang (num: N)' or similar for 'worker' or 'vector'
+   clauses) with the given CLAUSE_CODE.  Does nothing if neither PARENT_CLAUSE
+   nor LOOP_CLAUSE exist.  Returns the new clauses.  */
+
+static tree
+add_parent_or_loop_num_clause (tree parent_clause, tree loop_clause,
+			       omp_clause_code clause_code, tree clauses)
+{
+  if (parent_clause != NULL)
+    {
+      tree num_clause = unshare_expr (parent_clause);
+      OMP_CLAUSE_CHAIN (num_clause) = clauses;
+      clauses = num_clause;
+    }
+  else if (loop_clause != NULL)
+    {
+      /* The kernels region does not have a 'num_gangs' clause, but the loop
+	 itself had a 'gang (num: N)' clause.  Honor it by adding a
+	 'num_gangs (N)' clause on the compute construct.  */
+      tree num = OMP_CLAUSE_OPERAND (loop_clause, 0);
+      tree new_num_clause
+	= build_omp_clause (OMP_CLAUSE_LOCATION (loop_clause), clause_code);
+      OMP_CLAUSE_OPERAND (new_num_clause, 0) = num;
+      OMP_CLAUSE_CHAIN (new_num_clause) = clauses;
+      clauses = new_num_clause;
+    }
+  return clauses;
+}
+
+/* Helper for make_region_loop_nest, looking for 'worker (num: N)' or 'vector
+   (length: N)' clauses in nested loops.  Removes the argument, transferring it
+   to the enclosing compute construct (via WI->INFO).  If arguments within the
+   same loop nest conflict, emits a warning.
+
+   This function also decides whether to add an 'auto' clause on each of these
+   nested loops.  */
+
+struct adjust_nested_loop_clauses_wi_info
+{
+  tree *loop_gang_clause_ptr;
+  tree *loop_worker_clause_ptr;
+  tree *loop_vector_clause_ptr;
+};
+
+static tree
+adjust_nested_loop_clauses (gimple_stmt_iterator *gsi_p, bool *,
+			    struct walk_stmt_info *wi)
+{
+  struct adjust_nested_loop_clauses_wi_info *wi_info
+    = (struct adjust_nested_loop_clauses_wi_info *) wi->info;
+  gimple *stmt = gsi_stmt (*gsi_p);
+
+  if (gimple_code (stmt) == GIMPLE_OMP_FOR)
+    {
+      bool add_auto_clause = true;
+      tree loop_clauses = gimple_omp_for_clauses (stmt);
+      tree loop_clause = loop_clauses;
+      for (; loop_clause; loop_clause = OMP_CLAUSE_CHAIN (loop_clause))
+	{
+	  tree *outer_clause_ptr = NULL;
+	  switch (OMP_CLAUSE_CODE (loop_clause))
+	    {
+	    case OMP_CLAUSE_GANG:
+	      outer_clause_ptr = wi_info->loop_gang_clause_ptr;
+	      break;
+	    case OMP_CLAUSE_WORKER:
+	      outer_clause_ptr = wi_info->loop_worker_clause_ptr;
+	      break;
+	    case OMP_CLAUSE_VECTOR:
+	      outer_clause_ptr = wi_info->loop_vector_clause_ptr;
+	      break;
+	    case OMP_CLAUSE_SEQ:
+	    case OMP_CLAUSE_INDEPENDENT:
+	    case OMP_CLAUSE_AUTO:
+	      add_auto_clause = false;
+	    default:
+	      break;
+	    }
+	  if (outer_clause_ptr != NULL)
+	    {
+	      if (OMP_CLAUSE_OPERAND (loop_clause, 0) != NULL
+		  && *outer_clause_ptr == NULL)
+		{
+		  /* Transfer the clause to the enclosing compute construct and
+		     remove the numerical argument from the 'loop'.  */
+		  *outer_clause_ptr = unshare_expr (loop_clause);
+		  OMP_CLAUSE_OPERAND (loop_clause, 0) = NULL;
+		}
+	      else if (OMP_CLAUSE_OPERAND (loop_clause, 0) != NULL &&
+		       OMP_CLAUSE_OPERAND (*outer_clause_ptr, 0) != NULL)
+		{
+		  /* See if both of these are the same constant.  If they
+		     aren't, emit a warning.  */
+		  tree old_op = OMP_CLAUSE_OPERAND (*outer_clause_ptr, 0);
+		  tree new_op = OMP_CLAUSE_OPERAND (loop_clause, 0);
+		  if (!(cst_and_fits_in_hwi (old_op) &&
+			cst_and_fits_in_hwi (new_op) &&
+			int_cst_value (old_op) == int_cst_value (new_op)))
+		    {
+		      const char *clause_name
+			= omp_clause_code_name[OMP_CLAUSE_CODE (loop_clause)];
+		      error_at (gimple_location (stmt),
+				"cannot honor conflicting %qs clause",
+				clause_name);
+		      inform (OMP_CLAUSE_LOCATION (*outer_clause_ptr),
+			      "location of the previous clause"
+			      " in the same loop nest");
+		    }
+		  OMP_CLAUSE_OPERAND (loop_clause, 0) = NULL;
+		}
+	    }
+	}
+      if (add_auto_clause)
+	{
+	  tree auto_clause
+	    = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_AUTO);
+	  OMP_CLAUSE_CHAIN (auto_clause) = loop_clauses;
+	  gimple_omp_for_set_clauses (stmt, auto_clause);
+	}
+    }
+
+  return NULL;
+}
+
+/* Helper for make_region_loop_nest.  Transform OpenACC 'kernels'/'loop'
+   construct clauses into OpenACC 'parallel'/'loop' construct ones.  */
+
+static tree
+transform_kernels_loop_clauses (gimple *omp_for,
+				tree num_gangs_clause,
+				tree num_workers_clause,
+				tree vector_length_clause,
+				tree clauses)
+{
+  /* If this loop in a kernels region does not have an explicit 'seq',
+     'independent', or 'auto' clause, we must give it an explicit 'auto'
+     clause.
+     We also check for 'gang (num: N)' clauses.  These must not appear in
+     kernels regions that have their own 'num_gangs' clause.  Otherwise, they
+     must be converted and put on the region; similarly for 'worker' and
+     'vector' clauses.  */
+  bool add_auto_clause = true;
+  tree loop_gang_clause = NULL, loop_worker_clause = NULL,
+       loop_vector_clause = NULL;
+  tree loop_clauses = gimple_omp_for_clauses (omp_for);
+  for (tree loop_clause = loop_clauses;
+       loop_clause;
+       loop_clause = OMP_CLAUSE_CHAIN (loop_clause))
+    {
+      bool found_num_clause = false;
+      tree *clause_ptr, clause_to_check;
+      switch (OMP_CLAUSE_CODE (loop_clause))
+	{
+	case OMP_CLAUSE_GANG:
+	  found_num_clause = true;
+	  clause_ptr = &loop_gang_clause;
+	  clause_to_check = num_gangs_clause;
+	  break;
+	case OMP_CLAUSE_WORKER:
+	  found_num_clause = true;
+	  clause_ptr = &loop_worker_clause;
+	  clause_to_check = num_workers_clause;
+	  break;
+	case OMP_CLAUSE_VECTOR:
+	  found_num_clause = true;
+	  clause_ptr = &loop_vector_clause;
+	  clause_to_check = vector_length_clause;
+	  break;
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_AUTO:
+	  add_auto_clause = false;
+	default:
+	  break;
+	}
+      if (found_num_clause && OMP_CLAUSE_OPERAND (loop_clause, 0) != NULL)
+	{
+	  if (clause_to_check)
+	    {
+	      const char *clause_name
+		= omp_clause_code_name[OMP_CLAUSE_CODE (loop_clause)];
+	      const char *parent_clause_name
+		= omp_clause_code_name[OMP_CLAUSE_CODE (clause_to_check)];
+	      error_at (OMP_CLAUSE_LOCATION (loop_clause),
+			"argument not permitted on %qs clause"
+			" in OpenACC %<kernels%> region with a %qs clause",
+			clause_name, parent_clause_name);
+	      inform (OMP_CLAUSE_LOCATION (clause_to_check),
+		      "location of OpenACC %<kernels%>");
+	    }
+	  /* Copy the 'gang (N)'/'worker (N)'/'vector (N)' clause to the
+	     enclosing compute construct.  */
+	  *clause_ptr = unshare_expr (loop_clause);
+	  OMP_CLAUSE_CHAIN (*clause_ptr) = NULL;
+	  /* Leave a 'gang'/'worker'/'vector' clause on the 'loop', but without
+	     argument.  */
+	  OMP_CLAUSE_OPERAND (loop_clause, 0) = NULL;
+	}
+    }
+  if (add_auto_clause)
+    {
+      tree auto_clause = build_omp_clause (gimple_location (omp_for),
+					   OMP_CLAUSE_AUTO);
+      OMP_CLAUSE_CHAIN (auto_clause) = loop_clauses;
+      loop_clauses = auto_clause;
+    }
+  gimple_omp_for_set_clauses (omp_for, loop_clauses);
+  /* We must also recurse into the loop; it might contain nested loops having
+     their own 'worker (num: W)' or 'vector (length: V)' clauses.  Turn these
+     into 'worker'/'vector' clauses on the compute construct.  */
+  struct walk_stmt_info wi;
+  memset (&wi, 0, sizeof (wi));
+  struct adjust_nested_loop_clauses_wi_info wi_info;
+  wi_info.loop_gang_clause_ptr = &loop_gang_clause;
+  wi_info.loop_worker_clause_ptr = &loop_worker_clause;
+  wi_info.loop_vector_clause_ptr = &loop_vector_clause;
+  wi.info = &wi_info;
+  gimple *body = gimple_omp_body (omp_for);
+  walk_gimple_seq (body, adjust_nested_loop_clauses, NULL, &wi);
+  /* Check if there were conflicting numbers of workers or vector length.  */
+  if (loop_gang_clause != NULL &&
+      OMP_CLAUSE_OPERAND (loop_gang_clause, 0) == NULL)
+    loop_gang_clause = NULL;
+  if (loop_worker_clause != NULL &&
+      OMP_CLAUSE_OPERAND (loop_worker_clause, 0) == NULL)
+    loop_worker_clause = NULL;
+  if (loop_vector_clause != NULL &&
+      OMP_CLAUSE_OPERAND (loop_vector_clause, 0) == NULL)
+    vector_length_clause = NULL;
+
+  /* If the kernels region had 'num_gangs', 'num_worker', 'vector_length'
+     clauses, add these to this new compute construct.  */
+  clauses
+    = add_parent_or_loop_num_clause (num_gangs_clause, loop_gang_clause,
+				     OMP_CLAUSE_NUM_GANGS, clauses);
+  clauses
+    = add_parent_or_loop_num_clause (num_workers_clause, loop_worker_clause,
+				     OMP_CLAUSE_NUM_WORKERS, clauses);
+  clauses
+    = add_parent_or_loop_num_clause (vector_length_clause, loop_vector_clause,
+				     OMP_CLAUSE_VECTOR_LENGTH, clauses);
+
+  return clauses;
+}
+
+/* Construct a possibly gang-parallel compute construct containing the STMT,
+   which must be identical to, or a bind containing, the loop OMP_FOR.
+
+   The NUM_GANGS_CLAUSE, NUM_WORKERS_CLAUSE, and VECTOR_LENGTH_CLAUSE are
+   optional clauses from the original kernels region and must not be contained
+   in the other CLAUSES. The newly created compute construct is annotated with
+   the optional NUM_GANGS_CLAUSE as well as the other CLAUSES.  If there is no
+   NUM_GANGS_CLAUSE but the loop has a 'gang (num: N)' clause, that is
+   converted to a 'num_gangs (N)' clause on the new compute construct, and
+   similarly for 'worker' and 'vector' clauses.
+
+   The outermost loop gets an 'auto' clause unless there already is an
+   'seq'/'independent'/'auto' clause.  Nested loops inside OMP_FOR are treated
+   similarly by the adjust_nested_loop_clauses function.  */
+
+static gimple *
+make_region_loop_nest (gimple *omp_for, gimple_seq stmts,
+		       tree num_gangs_clause,
+		       tree num_workers_clause,
+		       tree vector_length_clause,
+		       tree clauses)
+{
+  /* This correctly unshares the entire clause chain rooted here.  */
+  clauses = unshare_expr (clauses);
+
+  /* Figure out the region code for this region.  */
+  /* Optimistic default: assume that the loop nest is parallelizable
+     (essentially, no GIMPLE_OMP_FOR with (explicit or implicit) 'auto' clause,
+     and no un-annotated loops).  */
+  int region_code = GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED;
+  adjust_region_code (stmts, &region_code);
+
+  if (region_code == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED)
+    {
+      if (dump_enabled_p ())
+	/* This is not MSG_OPTIMIZED_LOCATIONS, as we're just doing what the
+	   user asked us to.  */
+	dump_printf_loc (MSG_NOTE, omp_for,
+			 "parallelized loop nest"
+			 " in OpenACC %<kernels%> region\n");
+
+      clauses = transform_kernels_loop_clauses (omp_for,
+						num_gangs_clause,
+						num_workers_clause,
+						vector_length_clause,
+						clauses);
+    }
+  else if (region_code == GF_OMP_TARGET_KIND_OACC_KERNELS)
+    {
+      if (dump_enabled_p ())
+	dump_printf_loc (MSG_NOTE, omp_for,
+			 "forwarded loop nest"
+			 " in OpenACC %<kernels%> region"
+			 " to %<parloops%> for analysis\n");
+
+      /* We're transforming one 'GF_OMP_TARGET_KIND_OACC_KERNELS' into another
+	 'GF_OMP_TARGET_KIND_OACC_KERNELS', so don't have to
+	 'transform_kernels_loop_clauses'.  */
+      /* Re-assemble the clauses stripped off earlier.  */
+      clauses
+	= add_parent_or_loop_num_clause (num_gangs_clause, NULL,
+					 OMP_CLAUSE_NUM_GANGS, clauses);
+      clauses
+	= add_parent_or_loop_num_clause (num_workers_clause, NULL,
+					 OMP_CLAUSE_NUM_WORKERS, clauses);
+      clauses
+	= add_parent_or_loop_num_clause (vector_length_clause, NULL,
+					 OMP_CLAUSE_VECTOR_LENGTH, clauses);
+    }
+  else
+    gcc_unreachable ();
+
+  gimple *parallel_body_bind
+    = gimple_build_bind (NULL, stmts, make_node (BLOCK));
+  gimple *parallel_region
+    = gimple_build_omp_target (parallel_body_bind, region_code, clauses);
+  gimple_set_location (parallel_region, gimple_location (omp_for));
+
+  return parallel_region;
+}
+
+/* Eliminate any binds directly inside BIND by adding their statements to
+   BIND (i.e., modifying it in place), excluding binds that hold only an
+   OMP_FOR loop and associated setup/cleanup code.  Recurse into binds but
+   not other statements.  Return a chain of the local variables of eliminated
+   binds, i.e., the local variables found in nested binds.  If
+   INCLUDE_TOPLEVEL_VARS is true, this also includes the variables belonging
+   to BIND itself. */
+
+static tree
+flatten_binds (gbind *bind, bool include_toplevel_vars = false)
+{
+  tree vars = NULL, last_var = NULL;
+
+  if (include_toplevel_vars)
+    {
+      vars = gimple_bind_vars (bind);
+      last_var = vars;
+    }
+
+  gimple_seq new_body = NULL;
+  gimple_seq body_sequence = gimple_bind_body (bind);
+  gimple_stmt_iterator gsi, gsi_n;
+  for (gsi = gsi_start (body_sequence); !gsi_end_p (gsi); gsi = gsi_n)
+    {
+      /* Advance the iterator here because otherwise it would be invalidated
+	 by moving statements below.  */
+      gsi_n = gsi;
+      gsi_next (&gsi_n);
+
+      gimple *stmt = gsi_stmt (gsi);
+      /* Flatten bind statements, except the ones that contain only an
+	 OpenACC for loop.  */
+      if (gimple_code (stmt) == GIMPLE_BIND
+	  && !top_level_omp_for_in_stmt (stmt))
+	{
+	  gbind *inner_bind = as_a <gbind *> (stmt);
+	  /* Flatten recursively, and collect all variables.  */
+	  tree inner_vars = flatten_binds (inner_bind, true);
+	  gimple_seq inner_sequence = gimple_bind_body (inner_bind);
+	  gcc_assert (gimple_code (inner_sequence) != GIMPLE_BIND
+		      || top_level_omp_for_in_stmt (inner_sequence));
+	  gimple_seq_add_seq (&new_body, inner_sequence);
+	  /* Find the last variable; we will append others to it.  */
+	  while (last_var != NULL && TREE_CHAIN (last_var) != NULL)
+	    last_var = TREE_CHAIN (last_var);
+	  if (last_var != NULL)
+	    {
+	      TREE_CHAIN (last_var) = inner_vars;
+	      last_var = inner_vars;
+	    }
+	  else
+	    {
+	      vars = inner_vars;
+	      last_var = vars;
+	    }
+	}
+      else
+	gimple_seq_add_stmt (&new_body, stmt);
+    }
+
+  /* Put the possibly transformed body back into the bind.  */
+  gimple_bind_set_body (bind, new_body);
+  return vars;
+}
+
+/* Helper function for places where we construct data regions.  Wraps the BODY
+   inside a try-finally construct at LOC that calls __builtin_GOACC_data_end
+   in its cleanup block.  Returns this try statement.  */
+
+static gimple *
+make_data_region_try_statement (location_t loc, gimple *body)
+{
+  tree data_end_fn = builtin_decl_explicit (BUILT_IN_GOACC_DATA_END);
+  gimple *call = gimple_build_call (data_end_fn, 0);
+  gimple_seq cleanup = NULL;
+  gimple_seq_add_stmt (&cleanup, call);
+  gimple *try_stmt = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
+  gimple_set_location (body, loc);
+  return try_stmt;
+}
+
+/* If INNER_BIND_VARS holds variables, build an OpenACC data region with
+   location LOC containing BODY and having 'create (var)' clauses for each
+   variable.  If INNER_CLEANUP is present, add a try-finally statement with
+   this cleanup code in the finally block.  Return the new data region, or
+   the original BODY if no data region was needed.  */
+
+static gimple *
+maybe_build_inner_data_region (location_t loc, gimple *body,
+			       tree inner_bind_vars, gimple *inner_cleanup)
+{
+  /* Build data 'create (var)' clauses for these local variables.
+     Below we will add these to a data region enclosing the entire body
+     of the decomposed kernels region.  */
+  tree prev_mapped_var = NULL, next = NULL, artificial_vars = NULL,
+       inner_data_clauses = NULL;
+  for (tree v = inner_bind_vars; v; v = next)
+    {
+      next = TREE_CHAIN (v);
+      if (DECL_ARTIFICIAL (v)
+	  || TREE_CODE (v) == CONST_DECL
+	  || (DECL_LANG_SPECIFIC (current_function_decl)
+	      && DECL_TEMPLATE_INSTANTIATION (current_function_decl)))
+	{
+	  /* If this is an artificial temporary, it need not be mapped.  We
+	     move its declaration into the bind inside the data region.
+	     Also avoid mapping variables if we are inside a template
+	     instantiation; the code does not contain all the copies to
+	     temporaries that would make this legal.  */
+	  TREE_CHAIN (v) = artificial_vars;
+	  artificial_vars = v;
+	  if (prev_mapped_var != NULL)
+	    TREE_CHAIN (prev_mapped_var) = next;
+	  else
+	    inner_bind_vars = next;
+	}
+      else
+	{
+	  /* Otherwise, build the map clause.  */
+	  tree new_clause = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (new_clause, GOMP_MAP_ALLOC);
+	  OMP_CLAUSE_DECL (new_clause) = v;
+	  OMP_CLAUSE_SIZE (new_clause) = DECL_SIZE_UNIT (v);
+	  OMP_CLAUSE_CHAIN (new_clause) = inner_data_clauses;
+	  inner_data_clauses = new_clause;
+
+	  prev_mapped_var = v;
+	}
+    }
+
+  if (artificial_vars)
+    body = gimple_build_bind (artificial_vars, body, make_node (BLOCK));
+
+  /* If we determined above that there are variables that need to be created
+     on the device, construct a data region for them and wrap the body
+     inside that.  */
+  if (inner_data_clauses != NULL)
+    {
+      gcc_assert (inner_bind_vars != NULL);
+      gimple *inner_data_region
+	= gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DATA_KERNELS,
+				   inner_data_clauses);
+      gimple_set_location (inner_data_region, loc);
+      /* Make sure __builtin_GOACC_data_end is called at the end.  */
+      gimple *try_stmt = make_data_region_try_statement (loc, body);
+      gimple_omp_set_body (inner_data_region, try_stmt);
+      gimple *bind_body;
+      if (inner_cleanup != NULL)
+	/* Clobber all the inner variables that need to be clobbered.  */
+	bind_body = gimple_build_try (inner_data_region, inner_cleanup,
+				      GIMPLE_TRY_FINALLY);
+      else
+	bind_body = inner_data_region;
+      body = gimple_build_bind (inner_bind_vars, bind_body, make_node (BLOCK));
+    }
+
+  return body;
+}
+
+/* Helper function of decompose_kernels_region_body.  The statements in
+   REGION_BODY are expected to be decomposed parts; add an 'async' clause to
+   each.  Also add a 'wait' directive at the end of the sequence.  */
+
+static void
+add_async_clauses_and_wait (location_t loc, gimple_seq *region_body)
+{
+  tree default_async_queue
+    = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
+  for (gimple_stmt_iterator gsi = gsi_start (*region_body);
+       !gsi_end_p (gsi);
+       gsi_next (&gsi))
+    {
+      gimple *stmt = gsi_stmt (gsi);
+      tree target_clauses = gimple_omp_target_clauses (stmt);
+      tree new_async_clause = build_omp_clause (loc, OMP_CLAUSE_ASYNC);
+      OMP_CLAUSE_OPERAND (new_async_clause, 0) = default_async_queue;
+      OMP_CLAUSE_CHAIN (new_async_clause) = target_clauses;
+      target_clauses = new_async_clause;
+      gimple_omp_target_set_clauses (as_a <gomp_target *> (stmt),
+				     target_clauses);
+    }
+  /* A '#pragma acc wait' is just a call 'GOACC_wait (acc_async_sync, 0)'.  */
+  tree wait_fn = builtin_decl_explicit (BUILT_IN_GOACC_WAIT);
+  tree sync_arg = build_int_cst (integer_type_node, GOMP_ASYNC_SYNC);
+  gimple *wait_call = gimple_build_call (wait_fn, 2,
+					 sync_arg, integer_zero_node);
+  gimple_set_location (wait_call, loc);
+  gimple_seq_add_stmt (region_body, wait_call);
+}
+
+/* Auxiliary analysis of the body of a kernels region, to determine for each
+   OpenACC loop whether it is control-dependent (i.e., not necessarily
+   executed every time the kernels region is entered) or not.
+   We say that a loop is control-dependent if there is some cond, switch, or
+   goto statement that jumps over it, forwards or backwards.  For example,
+   if the loop is controlled by an if statement, then a jump to the true
+   block, the false block, or from one of those blocks to the control flow
+   join point will necessarily jump over the loop.
+   This analysis implements an ad-hoc union-find data structure classifying
+   statements into "control-flow regions" as follows: Most statements are in
+   the same region as their predecessor, except that each OpenACC loop is in
+   a region of its own, and each OpenACC loop's successor starts a new
+   region.  We then unite the regions of any statements linked by jumps,
+   placing any cond, switch, or goto statement in the same region as its
+   target label(s).
+   In the end, control dependence of OpenACC loops can be determined by
+   comparing their immediate predecessor and successor statements' regions.
+   A jump crosses the loop if and only if the predecessor and successor are
+   in the same region.  (If there is no predecessor or successor, the loop
+   is executed unconditionally.)
+   The methods in this class identify statements by their index in the
+   kernels region's body.  */
+
+class control_flow_regions
+{
+  public:
+    /* Initialize an instance and pre-compute the control-flow region
+       information for the statement sequence SEQ.  */
+    control_flow_regions (gimple_seq seq);
+
+    /* Return true if the statement with the given index IDX in the analyzed
+       statement sequence is an unconditionally executed OpenACC loop.  */
+    bool is_unconditional_oacc_for_loop (size_t idx);
+
+  private:
+    /* Find the region representative for the statement identified by index
+       STMT_IDX.  */
+    size_t find_rep (size_t stmt_idx);
+
+    /* Union the regions containing the statements represented by
+       representatives A and B.  */
+    void union_reps (size_t a, size_t b);
+
+    /* Helper for the constructor.  Performs the actual computation of the
+       control-flow regions in the statement sequence SEQ.  */
+    void compute_regions (gimple_seq seq);
+
+    /* The mapping from statement indices to region representatives.  */
+    vec <size_t> representatives;
+
+    /* A cache mapping statement indices to a flag indicating whether the
+       statement is a top level OpenACC for loop.  */
+    vec <bool> omp_for_loops;
+};
+
+control_flow_regions::control_flow_regions (gimple_seq seq)
+{
+  representatives.create (1);
+  omp_for_loops.create (1);
+  compute_regions (seq);
+}
+
+bool
+control_flow_regions::is_unconditional_oacc_for_loop (size_t idx)
+{
+  if (idx == 0 || idx == representatives.length () - 1)
+    /* The first or last statement in the kernels region.  This means that
+       there is no room before or after it for a jump or a label.  Thus
+       there cannot be a jump across it, so it is unconditional.  */
+    return true;
+  /* Otherwise, the loop is unconditional if the statements before and after
+     it are in different control flow regions.  Scan forward and backward,
+     skipping over neighboring OpenACC for loops, to find these preceding
+     statements.  */
+  size_t prev_index = idx - 1;
+  while (prev_index > 0 && omp_for_loops [prev_index] == true)
+    prev_index--;
+  /* If all preceding statements are also OpenACC loops, all of these are
+     unconditional.  */
+  if (prev_index == 0)
+    return true;
+  size_t succ_index = idx + 1;
+  while (succ_index < omp_for_loops.length ()
+	 && omp_for_loops [succ_index] == true)
+    succ_index++;
+  /* If all following statements are also OpenACC loops, all of these are
+     unconditional.  */
+  if (succ_index == omp_for_loops.length ())
+    return true;
+  return (find_rep (prev_index) != find_rep (succ_index));
+}
+
+size_t
+control_flow_regions::find_rep (size_t stmt_idx)
+{
+  size_t rep = stmt_idx, aux = stmt_idx;
+  /* Find the root representative of this statement.  */
+  while (representatives[rep] != rep)
+    rep = representatives[rep];
+  /* Compress the path from the original statement to the representative.  */
+  while (representatives[aux] != rep)
+    {
+      size_t tmp = representatives[aux];
+      representatives[aux] = rep;
+      aux = tmp;
+    }
+  return rep;
+}
+
+void
+control_flow_regions::union_reps (size_t a, size_t b)
+{
+  a = find_rep (a);
+  b = find_rep (b);
+  representatives[b] = a;
+}
+
+void
+control_flow_regions::compute_regions (gimple_seq seq)
+{
+  hash_map <gimple *, size_t> control_flow_reps;
+  hash_map <tree, size_t> label_reps;
+  size_t current_region = 0, idx = 0;
+
+  /* In a first pass, assign an initial region to each statement.  Except in
+     the case of OpenACC loops, each statement simply gets the same region
+     representative as its predecessor.  */
+  for (gimple_stmt_iterator gsi = gsi_start (seq);
+       !gsi_end_p (gsi);
+       gsi_next (&gsi))
+    {
+      gimple *stmt = gsi_stmt (gsi);
+      gimple *omp_for = top_level_omp_for_in_stmt (stmt);
+      omp_for_loops.safe_push (omp_for != NULL);
+      if (omp_for != NULL)
+	{
+	  /* Assign a new region to this loop and to its successor.  */
+	  current_region = idx;
+	  representatives.safe_push (current_region);
+	  current_region++;
+	}
+      else
+	{
+	  representatives.safe_push (current_region);
+	  /* Remember any jumps and labels for the second pass below.  */
+	  if (gimple_code (stmt) == GIMPLE_COND
+	      || gimple_code (stmt) == GIMPLE_SWITCH
+	      || gimple_code (stmt) == GIMPLE_GOTO)
+	    control_flow_reps.put (stmt, current_region);
+	  else if (gimple_code (stmt) == GIMPLE_LABEL)
+	    label_reps.put (gimple_label_label (as_a <glabel *> (stmt)),
+			    current_region);
+	}
+      idx++;
+    }
+  gcc_assert (representatives.length () == omp_for_loops.length ());
+
+  /* Revisit all the control flow statements and union the region of each
+     cond, switch, or goto statement with the target labels' regions.  */
+  for (hash_map <gimple *, size_t>::iterator it = control_flow_reps.begin ();
+       it != control_flow_reps.end ();
+       ++it)
+    {
+      gimple *stmt = (*it).first;
+      size_t stmt_rep = (*it).second;
+      switch (gimple_code (stmt))
+	{
+	  tree label;
+	  unsigned int n;
+
+	case GIMPLE_COND:
+	  label = gimple_cond_true_label (as_a <gcond *> (stmt));
+	  union_reps (stmt_rep, *label_reps.get (label));
+	  label = gimple_cond_false_label (as_a <gcond *> (stmt));
+	  union_reps (stmt_rep, *label_reps.get (label));
+	  break;
+
+	case GIMPLE_SWITCH:
+	  n = gimple_switch_num_labels (as_a <gswitch *> (stmt));
+	  for (unsigned int i = 0; i < n; i++)
+	    {
+	      tree switch_case
+		= gimple_switch_label (as_a <gswitch *> (stmt), i);
+	      label = CASE_LABEL (switch_case);
+	      union_reps (stmt_rep, *label_reps.get (label));
+	    }
+	  break;
+
+	case GIMPLE_GOTO:
+	  label = gimple_goto_dest (stmt);
+	  union_reps (stmt_rep, *label_reps.get (label));
+	  break;
+
+	default:
+	  gcc_unreachable ();
+	}
+    }
+}
+
+/* Decompose the body of the KERNELS_REGION, which was originally annotated
+   with the KERNELS_CLAUSES, into a series of compute constructs.  */
+
+static gimple *
+decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
+{
+  location_t loc = gimple_location (kernels_region);
+
+  /* The kernels clauses will be propagated to the child clauses unmodified,
+     except that the 'num_gangs', 'num_workers', and 'vector_length' clauses
+     will only be added to loop regions.  The other regions are "gang-single"
+     and get an explicit 'num_gangs (1)' clause.  So separate out the
+     'num_gangs', 'num_workers', and 'vector_length' clauses here.
+     Also check for the presence of an 'async' clause but do not remove it from
+     the 'kernels' clauses.  */
+  tree num_gangs_clause = NULL, num_workers_clause = NULL,
+       vector_length_clause = NULL;
+  tree async_clause = NULL;
+  tree prev_clause = NULL, next_clause = NULL;
+  tree parallel_clauses = kernels_clauses;
+  for (tree c = parallel_clauses; c; c = next_clause)
+    {
+      /* Preserve this here, as we might NULL it later.  */
+      next_clause = OMP_CLAUSE_CHAIN (c);
+
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_GANGS
+	  || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_WORKERS
+	  || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR_LENGTH)
+	{
+	  /* Cut this clause out of the chain.  */
+	  if (prev_clause != NULL)
+	    OMP_CLAUSE_CHAIN (prev_clause) = OMP_CLAUSE_CHAIN (c);
+	  else
+	    kernels_clauses = OMP_CLAUSE_CHAIN (c);
+	  OMP_CLAUSE_CHAIN (c) = NULL;
+	  switch (OMP_CLAUSE_CODE (c))
+	    {
+	    case OMP_CLAUSE_NUM_GANGS:
+	      num_gangs_clause = c;
+	      break;
+	    case OMP_CLAUSE_NUM_WORKERS:
+	      num_workers_clause = c;
+	      break;
+	    case OMP_CLAUSE_VECTOR_LENGTH:
+	      vector_length_clause = c;
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	}
+      else
+	prev_clause = c;
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
+	async_clause = c;
+    }
+
+  gimple *kernels_body = gimple_omp_body (kernels_region);
+  gbind *kernels_bind = as_a <gbind *> (kernels_body);
+
+  /* The body of the region may contain other nested binds declaring inner
+     local variables.  Collapse all these binds into one to ensure that we
+     have a single sequence of statements to iterate over; also, collect all
+     inner variables.  */
+  tree inner_bind_vars = flatten_binds (kernels_bind);
+  gimple_seq body_sequence = gimple_bind_body (kernels_bind);
+
+  /* All these inner variables will get allocated on the device (below, by
+     calling maybe_build_inner_data_region).  Here we create 'present'
+     clauses for them and add these clauses to the list of clauses to be
+     attached to each inner compute construct.  */
+  tree present_clauses = kernels_clauses;
+  for (tree var = inner_bind_vars; var; var = TREE_CHAIN (var))
+    {
+      if (!DECL_ARTIFICIAL (var) && TREE_CODE (var) != CONST_DECL)
+	{
+	  tree present_clause = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (present_clause, GOMP_MAP_FORCE_PRESENT);
+	  OMP_CLAUSE_DECL (present_clause) = var;
+	  OMP_CLAUSE_SIZE (present_clause) = DECL_SIZE_UNIT (var);
+	  OMP_CLAUSE_CHAIN (present_clause) = present_clauses;
+	  present_clauses = present_clause;
+	}
+    }
+  kernels_clauses = present_clauses;
+
+  /* In addition to nested binds, the "real" body of the region may be
+     nested inside a try-finally block.  Find its cleanup block, which
+     contains code to clobber the local variables that must be clobbered.  */
+  gimple *inner_cleanup = NULL;
+  if (body_sequence != NULL && gimple_code (body_sequence) == GIMPLE_TRY)
+    {
+      if (gimple_seq_singleton_p (body_sequence))
+	{
+	  /* The try statement is the only thing inside the bind.  */
+	  inner_cleanup = gimple_try_cleanup (body_sequence);
+	  body_sequence = gimple_try_eval (body_sequence);
+	}
+      else
+	{
+	  /* The bind's body starts with a try statement, but it is followed
+	     by other things.  */
+	  gimple_stmt_iterator gsi = gsi_start (body_sequence);
+	  gimple *try_stmt = gsi_stmt (gsi);
+	  inner_cleanup = gimple_try_cleanup (try_stmt);
+	  gimple *try_body = gimple_try_eval (try_stmt);
+
+	  gsi_remove (&gsi, false);
+	  /* Now gsi indicates the sequence of statements after the try
+	     statement in the bind.  Append the statement in the try body and
+	     the trailing statements from gsi.  */
+	  gsi_insert_seq_before (&gsi, try_body, GSI_CONTINUE_LINKING);
+	  body_sequence = gsi_stmt (gsi);
+	}
+    }
+
+  /* This sequence will collect all the top-level statements in the body of
+     the data region we are about to construct.  */
+  gimple_seq region_body = NULL;
+  /* This sequence will collect consecutive statements to be put into a
+     gang-single region.  */
+  gimple_seq gang_single_seq = NULL;
+  /* Flag recording whether the gang_single_seq only contains copies to
+     local variables.  These may be loop setup code that should not be
+     separated from the loop.  */
+  bool only_simple_assignments = true;
+
+  /* Precompute the control flow region information to determine whether an
+     OpenACC loop is executed conditionally or unconditionally.  */
+  control_flow_regions cf_regions (body_sequence);
+
+  /* Iterate over the statements in the kernels region's body.  */
+  size_t idx = 0;
+  gimple_stmt_iterator gsi, gsi_n;
+  for (gsi = gsi_start (body_sequence); !gsi_end_p (gsi); gsi = gsi_n, idx++)
+    {
+      /* Advance the iterator here because otherwise it would be invalidated
+	 by moving statements below.  */
+      gsi_n = gsi;
+      gsi_next (&gsi_n);
+
+      gimple *stmt = gsi_stmt (gsi);
+      gimple *omp_for = top_level_omp_for_in_stmt (stmt);
+      bool is_unconditional_oacc_for_loop = false;
+      if (omp_for != NULL)
+	is_unconditional_oacc_for_loop
+	  = cf_regions.is_unconditional_oacc_for_loop (idx);
+      if (omp_for != NULL
+	  && is_unconditional_oacc_for_loop)
+	{
+	  /* This is an OMP for statement, put it into a separate region.
+	     But first, construct a gang-single region containing any
+	     complex sequential statements we may have seen.  */
+	  if (gang_single_seq != NULL && !only_simple_assignments)
+	    {
+	      gimple *single_region
+		= make_region_seq (loc, gang_single_seq,
+				   num_gangs_clause,
+				   num_workers_clause,
+				   vector_length_clause,
+				   kernels_clauses);
+	      gimple_seq_add_stmt (&region_body, single_region);
+	    }
+	  else if (gang_single_seq != NULL && only_simple_assignments)
+	    {
+	      /* There is a sequence of sequential statements preceding this
+		 loop, but they are all simple assignments.  This is
+		 probably setup code for the loop; in particular, Fortran DO
+		 loops are preceded by code to copy the loop limit variable
+		 to a temporary.  Group this code together with the loop
+		 itself.  */
+	      gimple_seq_add_stmt (&gang_single_seq, stmt);
+	      stmt = gimple_build_bind (NULL, gang_single_seq,
+					make_node (BLOCK));
+	    }
+	  gang_single_seq = NULL;
+	  only_simple_assignments = true;
+
+	  gimple_seq parallel_seq = NULL;
+	  gimple_seq_add_stmt (&parallel_seq, stmt);
+	  gimple *parallel_region
+	    = make_region_loop_nest (omp_for, parallel_seq,
+				     num_gangs_clause,
+				     num_workers_clause,
+				     vector_length_clause,
+				     kernels_clauses);
+	  gimple_seq_add_stmt (&region_body, parallel_region);
+	}
+      else
+	{
+	  if (omp_for != NULL)
+	    {
+	      gcc_checking_assert (!is_unconditional_oacc_for_loop);
+	      if (dump_enabled_p ())
+		dump_printf_loc (MSG_MISSED_OPTIMIZATION, omp_for,
+				 "unparallelized loop nest"
+				 " in OpenACC %<kernels%> region:"
+				 " it's executed conditionally\n");
+	    }
+
+	  /* This is not an unconditional OMP for statement, so it will be
+	     put into a gang-single region.  */
+	  gimple_seq_add_stmt (&gang_single_seq, stmt);
+	  /* Is this a simple assignment? We call it simple if it is an
+	     assignment to an artificial local variable.  This captures
+	     Fortran loop setup code computing loop bounds and offsets.  */
+	  bool is_simple_assignment
+	    = (gimple_code (stmt) == GIMPLE_ASSIGN
+	       && TREE_CODE (gimple_assign_lhs (stmt)) == VAR_DECL
+	       && DECL_ARTIFICIAL (gimple_assign_lhs (stmt)));
+	  if (!is_simple_assignment)
+	    only_simple_assignments = false;
+	}
+    }
+
+  /* If we did not emit a new region, and are not going to emit one now
+     (that is, the original region was empty), prepare to emit a dummy so as
+     to preserve the original construct, which other processing (at least
+     test cases) depend on.  */
+  if (region_body == NULL && gang_single_seq == NULL)
+    {
+      gimple *stmt = gimple_build_nop ();
+      gimple_set_location (stmt, loc);
+      gimple_seq_add_stmt (&gang_single_seq, stmt);
+    }
+
+  /* Gather up any remaining gang-single statements.  */
+  if (gang_single_seq != NULL)
+    {
+      gimple *single_region
+	= make_region_seq (loc, gang_single_seq,
+			   num_gangs_clause,
+			   num_workers_clause,
+			   vector_length_clause,
+			   kernels_clauses);
+      gimple_seq_add_stmt (&region_body, single_region);
+    }
+
+  /* We want to launch these kernels asynchronously.  If the original
+     kernels region had an async clause, this is done automatically because
+     that async clause was copied to the individual regions we created.
+     Otherwise, add an async clause to each newly created region, as well as
+     a wait directive at the end.  */
+  if (async_clause == NULL)
+    add_async_clauses_and_wait (loc, &region_body);
+
+  tree kernels_locals = gimple_bind_vars (as_a <gbind *> (kernels_body));
+  gimple *body = gimple_build_bind (kernels_locals, region_body,
+				    make_node (BLOCK));
+
+  /* If we found variables declared in nested scopes, build a data region to
+     map them to the device.  */
+  body = maybe_build_inner_data_region (loc, body, inner_bind_vars,
+					inner_cleanup);
+
+  return body;
+}
+
+/* Decompose one OpenACC 'kernels' construct into an OpenACC 'data' construct
+   containing the original OpenACC 'kernels' construct's region cut up into a
+   sequence of compute constructs.  */
+
+static gimple *
+omp_oacc_kernels_decompose_1 (gimple *kernels_stmt)
+{
+  gcc_checking_assert (gimple_omp_target_kind (kernels_stmt)
+		       == GF_OMP_TARGET_KIND_OACC_KERNELS);
+  location_t loc = gimple_location (kernels_stmt);
+
+  /* Collect the data clauses of the OpenACC 'kernels' directive and create a
+     new OpenACC 'data' construct with those clauses.  */
+  tree kernels_clauses = gimple_omp_target_clauses (kernels_stmt);
+  tree data_clauses = NULL;
+  for (tree c = kernels_clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      /* Certain clauses are copied to the enclosing OpenACC 'data'.  Other
+	 clauses remain on the OpenACC 'kernels'.  */
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	{
+	  tree decl = OMP_CLAUSE_DECL (c);
+	  HOST_WIDE_INT map_kind = OMP_CLAUSE_MAP_KIND (c);
+	  switch (map_kind)
+	    {
+	    default:
+	      if (map_kind == GOMP_MAP_ALLOC
+		  && integer_zerop (OMP_CLAUSE_SIZE (c)))
+		/* ??? This is an alloc clause for mapping a pointer whose
+		   target is already mapped.  We leave these on the inner
+		   compute constructs because moving them to the outer data
+		   region causes runtime errors.  */
+		break;
+
+	      /* For non-artificial variables, and for non-declaration
+		 expressions like A[0:n], copy the clause to the data
+		 region.  */
+	      if ((DECL_P (decl) && !DECL_ARTIFICIAL (decl))
+		  || !DECL_P (decl))
+		{
+		  tree new_clause = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						      OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND (new_clause, map_kind);
+		  /* This must be unshared here to avoid "incorrect sharing
+		     of tree nodes" errors from verify_gimple.  */
+		  OMP_CLAUSE_DECL (new_clause) = unshare_expr (decl);
+		  OMP_CLAUSE_SIZE (new_clause) = OMP_CLAUSE_SIZE (c);
+		  OMP_CLAUSE_CHAIN (new_clause) = data_clauses;
+		  data_clauses = new_clause;
+
+		  /* Now that this data is mapped, turn the data clause on the
+		     inner OpenACC 'kernels' into a 'present' clause.  */
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT);
+		}
+	      break;
+
+	    case GOMP_MAP_POINTER:
+	    case GOMP_MAP_TO_PSET:
+	    case GOMP_MAP_FORCE_TOFROM:
+	    case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+	      /* ??? Copying these map kinds leads to internal compiler
+		 errors in later passes.  */
+	      break;
+	    }
+	}
+      else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF)
+	{
+	  /* If there is an 'if' clause, it must be duplicated to the
+	     enclosing data region.  Temporarily remove the if clause's
+	     chain to avoid copying it.  */
+	  tree saved_chain = OMP_CLAUSE_CHAIN (c);
+	  OMP_CLAUSE_CHAIN (c) = NULL;
+	  tree new_if_clause = unshare_expr (c);
+	  OMP_CLAUSE_CHAIN (c) = saved_chain;
+	  OMP_CLAUSE_CHAIN (new_if_clause) = data_clauses;
+	  data_clauses = new_if_clause;
+	}
+    }
+  /* Restore the original order of the clauses.  */
+  data_clauses = nreverse (data_clauses);
+
+  gimple *data_region
+    = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DATA_KERNELS,
+			       data_clauses);
+  gimple_set_location (data_region, loc);
+
+  /* Transform the body of the kernels region into a sequence of compute
+     constructs.  */
+  gimple *body = decompose_kernels_region_body (kernels_stmt,
+						kernels_clauses);
+
+  /* Put the transformed pieces together.  The entire body of the region is
+     wrapped in a try-finally statement that calls __builtin_GOACC_data_end
+     for cleanup.  */
+  gimple *try_stmt = make_data_region_try_statement (loc, body);
+  gimple_omp_set_body (data_region, try_stmt);
+
+  return data_region;
+}
+
+
+/* Decompose OpenACC 'kernels' constructs in the current function.  */
+
+static tree
+omp_oacc_kernels_decompose_callback_stmt (gimple_stmt_iterator *gsi_p,
+					  bool *handled_ops_p,
+					  struct walk_stmt_info *)
+{
+  gimple *stmt = gsi_stmt (*gsi_p);
+
+  if ((gimple_code (stmt) == GIMPLE_OMP_TARGET)
+      && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS)
+    {
+      gimple *stmt_new = omp_oacc_kernels_decompose_1 (stmt);
+      gsi_replace (gsi_p, stmt_new, false);
+      *handled_ops_p = true;
+    }
+  else
+    *handled_ops_p = false;
+
+  return NULL;
+}
+
+static unsigned int
+omp_oacc_kernels_decompose (void)
+{
+  gimple_seq body = gimple_body (current_function_decl);
+
+  struct walk_stmt_info wi;
+  memset (&wi, 0, sizeof (wi));
+  walk_gimple_seq_mod (&body, omp_oacc_kernels_decompose_callback_stmt, NULL,
+		       &wi);
+
+  gimple_set_body (current_function_decl, body);
+
+  return 0;
+}
+
+
+namespace {
+
+const pass_data pass_data_omp_oacc_kernels_decompose =
+{
+  GIMPLE_PASS, /* type */
+  "omp_oacc_kernels_decompose", /* name */
+  OPTGROUP_OMP, /* optinfo_flags */
+  TV_NONE, /* tv_id */
+  PROP_gimple_any, /* properties_required */
+  0, /* properties_provided */
+  0, /* properties_destroyed */
+  0, /* todo_flags_start */
+  0, /* todo_flags_finish */
+};
+
+class pass_omp_oacc_kernels_decompose : public gimple_opt_pass
+{
+public:
+  pass_omp_oacc_kernels_decompose (gcc::context *ctxt)
+    : gimple_opt_pass (pass_data_omp_oacc_kernels_decompose, ctxt)
+  {}
+
+  /* opt_pass methods: */
+  virtual bool gate (function *)
+  {
+    return (flag_openacc
+	    && flag_openacc_kernels == OPENACC_KERNELS_DECOMPOSE);
+  }
+  virtual unsigned int execute (function *)
+  {
+    return omp_oacc_kernels_decompose ();
+  }
+
+}; // class pass_omp_oacc_kernels_decompose
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_omp_oacc_kernels_decompose (gcc::context *ctxt)
+{
+  return new pass_omp_oacc_kernels_decompose (ctxt);
+}