diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 3d9b06d..79030d6 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,39 @@
+2014-02-28  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gimple.def (GIMPLE_OACC_KERNELS): New code.
+	* doc/gimple.texi: Document it.
+	* gimple.h (gimple_has_substatements, CASE_GIMPLE_OMP)
+	(is_gimple_omp_oacc_specifically): Handle it.
+	(gimple_statement_oacc_kernels): New struct.
+	(gimple_build_oacc_kernels): New prototype.
+	(gimple_oacc_kernels_clauses, gimple_oacc_kernels_clauses_ptr)
+	(gimple_oacc_kernels_set_clauses, gimple_oacc_kernels_child_fn)
+	(gimple_oacc_kernels_child_fn_ptr)
+	(gimple_oacc_kernels_set_child_fn, gimple_oacc_kernels_data_arg)
+	(gimple_oacc_kernels_data_arg_ptr)
+	(gimple_oacc_kernels_set_data_arg): New inline functions.
+	* gimple.c (gimple_build_oacc_kernels): New function.
+	(gimple_copy): Handle GIMPLE_OACC_KERNELS.
+	* gimple-low.c (lower_stmt): Likewise.
+	* gimple-walk.c (walk_gimple_op, walk_gimple_stmt): Likewise.
+	* gimple-pretty-print.c (pp_gimple_stmt_1): Likewise.
+	(dump_gimple_oacc_parallel): Rename to dump_gimple_oacc_offload.
+	Also handle GIMPLE_OACC_KERNELS.  Update all callers.
+	* gimplify.c (gimplify_omp_workshare, gimplify_expr): Handle
+	OACC_KERNELS.
+	* oacc-builtins.def (BUILT_IN_GOACC_KERNELS): New builtin.
+	* omp-low.c (scan_oacc_parallel, expand_oacc_parallel)
+	(lower_oacc_parallel): Rename to scan_oacc_offload,
+	expand_oacc_offload, and lower_oacc_offload.  Also handle
+	GIMPLE_OACC_KERNELS.  Update all callers.
+	(scan_sharing_clauses, scan_omp_1_stmt, expand_omp, lower_omp_1)
+	(diagnose_sb_0, diagnose_sb_1, diagnose_sb_2)
+	(make_gimple_omp_edges): Handle GIMPLE_OACC_KERNELS.
+	* tree-inline.c (remap_gimple_stmt, estimate_num_insns): Likewise.
+	* tree-nested.c (convert_nonlocal_reference_stmt)
+	(convert_local_reference_stmt, convert_tramp_reference_stmt)
+	(convert_gimple_call): Likewise.
+
 2014-02-27  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* gimplify.c (gimplify_oacc_parallel): Merge into
diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi
index 6136963..91748a6 100644
--- gcc/doc/gimple.texi
+++ gcc/doc/gimple.texi
@@ -338,6 +338,7 @@ The following table briefly describes the GIMPLE instruction set.
 @item @code{GIMPLE_GOTO}		@tab x			@tab x
 @item @code{GIMPLE_LABEL}		@tab x			@tab x
 @item @code{GIMPLE_NOP}			@tab x			@tab x
+@item @code{GIMPLE_OACC_KERNELS}	@tab x			@tab x
 @item @code{GIMPLE_OACC_PARALLEL}	@tab x			@tab x
 @item @code{GIMPLE_OMP_ATOMIC_LOAD}	@tab x			@tab x
 @item @code{GIMPLE_OMP_ATOMIC_STORE}	@tab x			@tab x
@@ -906,6 +907,7 @@ Return a deep copy of statement @code{STMT}.
 * @code{GIMPLE_EH_FILTER}::
 * @code{GIMPLE_LABEL}::
 * @code{GIMPLE_NOP}::
+* @code{GIMPLE_OACC_KERNELS}::
 * @code{GIMPLE_OACC_PARALLEL}::
 * @code{GIMPLE_OMP_ATOMIC_LOAD}::
 * @code{GIMPLE_OMP_ATOMIC_STORE}::
@@ -1553,6 +1555,11 @@ Returns @code{TRUE} if statement @code{G} is a @code{GIMPLE_NOP}.
 @end deftypefn
 
 
+@node @code{GIMPLE_OACC_KERNELS}
+@subsection @code{GIMPLE_OACC_KERNELS}
+@cindex @code{GIMPLE_OACC_KERNELS}
+
+
 @node @code{GIMPLE_OACC_PARALLEL}
 @subsection @code{GIMPLE_OACC_PARALLEL}
 @cindex @code{GIMPLE_OACC_PARALLEL}
diff --git gcc/gimple-low.c gcc/gimple-low.c
index 7bf69bd..c7d9c1c 100644
--- gcc/gimple-low.c
+++ gcc/gimple-low.c
@@ -353,6 +353,7 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
       }
       break;
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index ad9369c..1a31192 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1843,36 +1843,57 @@ dump_gimple_phi (pretty_printer *buffer, gimple phi, int spc, bool comment,
 }
 
 
-/* Dump a GIMPLE_OACC_PARALLEL tuple on the pretty_printer BUFFER, SPC spaces
+/* Dump an OpenACC offload tuple on the pretty_printer BUFFER, SPC spaces
    of indent.  FLAGS specifies details to show in the dump (see TDF_* in
    dumpfile.h).  */
 
 static void
-dump_gimple_oacc_parallel (pretty_printer *buffer, gimple gs, int spc,
-                          int flags)
+dump_gimple_oacc_offload (pretty_printer *buffer, gimple gs, int spc,
+			  int flags)
 {
+  tree (*gimple_omp_clauses) (const_gimple);
+  tree (*gimple_omp_child_fn) (const_gimple);
+  tree (*gimple_omp_data_arg) (const_gimple);
+  const char *kind;
+  switch (gimple_code (gs))
+    {
+    case GIMPLE_OACC_KERNELS:
+      gimple_omp_clauses = gimple_oacc_kernels_clauses;
+      gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
+      gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+      kind = "kernels";
+      break;
+    case GIMPLE_OACC_PARALLEL:
+      gimple_omp_clauses = gimple_oacc_parallel_clauses;
+      gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
+      gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
+      kind = "parallel";
+      break;
+    default:
+      gcc_unreachable ();
+    }
   if (flags & TDF_RAW)
     {
       dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs,
                        gimple_omp_body (gs));
-      dump_omp_clauses (buffer, gimple_oacc_parallel_clauses (gs), spc, flags);
+      dump_omp_clauses (buffer, gimple_omp_clauses (gs), spc, flags);
       dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>",
-                       gimple_oacc_parallel_child_fn (gs),
-                       gimple_oacc_parallel_data_arg (gs));
+                       gimple_omp_child_fn (gs), gimple_omp_data_arg (gs));
     }
   else
     {
       gimple_seq body;
-      pp_string (buffer, "#pragma acc parallel");
-      dump_omp_clauses (buffer, gimple_oacc_parallel_clauses (gs), spc, flags);
-      if (gimple_oacc_parallel_child_fn (gs))
+      pp_string (buffer, "#pragma acc ");
+      pp_string (buffer, kind);
+      dump_omp_clauses (buffer, gimple_omp_clauses (gs), spc, flags);
+      if (gimple_omp_child_fn (gs))
 	{
 	  pp_string (buffer, " [child fn: ");
-	  dump_generic_node (buffer, gimple_oacc_parallel_child_fn (gs),
+	  dump_generic_node (buffer, gimple_omp_child_fn (gs),
 			     spc, flags, false);
 	  pp_string (buffer, " (");
-	  if (gimple_oacc_parallel_data_arg (gs))
-	    dump_generic_node (buffer, gimple_oacc_parallel_data_arg (gs),
+	  if (gimple_omp_data_arg (gs))
+	    dump_generic_node (buffer, gimple_omp_data_arg (gs),
 			       spc, flags, false);
 	  else
 	    pp_string (buffer, "???");
@@ -2193,8 +2214,9 @@ pp_gimple_stmt_1 (pretty_printer *buffer, gimple gs, int spc, int flags)
       dump_gimple_phi (buffer, gs, spc, false, flags);
       break;
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
-      dump_gimple_oacc_parallel (buffer, gs, spc, flags);
+      dump_gimple_oacc_offload (buffer, gs, spc, flags);
       break;
 
     case GIMPLE_OMP_PARALLEL:
diff --git gcc/gimple-walk.c gcc/gimple-walk.c
index a90ba55..b5b4095 100644
--- gcc/gimple-walk.c
+++ gcc/gimple-walk.c
@@ -296,6 +296,21 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op,
 	return ret;
       break;
 
+    case GIMPLE_OACC_KERNELS:
+      ret = walk_tree (gimple_oacc_kernels_clauses_ptr (stmt), callback_op,
+		       wi, pset);
+      if (ret)
+	return ret;
+      ret = walk_tree (gimple_oacc_kernels_child_fn_ptr (stmt), callback_op,
+		       wi, pset);
+      if (ret)
+	return ret;
+      ret = walk_tree (gimple_oacc_kernels_data_arg_ptr (stmt), callback_op,
+		       wi, pset);
+      if (ret)
+	return ret;
+      break;
+
     case GIMPLE_OACC_PARALLEL:
       ret = walk_tree (gimple_oacc_parallel_clauses_ptr (stmt), callback_op,
 		       wi, pset);
@@ -606,6 +621,7 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
 	return wi->callback_result;
 
       /* FALL THROUGH.  */
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_CRITICAL:
     case GIMPLE_OMP_MASTER:
diff --git gcc/gimple.c gcc/gimple.c
index 30561b1..1862de2 100644
--- gcc/gimple.c
+++ gcc/gimple.c
@@ -799,6 +799,23 @@ gimple_build_debug_source_bind_stat (tree var, tree value,
 }
 
 
+/* Build a GIMPLE_OACC_KERNELS statement.
+
+   BODY is sequence of statements which are executed as kernels.
+   CLAUSES are the OpenACC kernels construct's clauses.  */
+
+gimple
+gimple_build_oacc_kernels (gimple_seq body, tree clauses)
+{
+  gimple p = gimple_alloc (GIMPLE_OACC_KERNELS, 0);
+  if (body)
+    gimple_omp_set_body (p, body);
+  gimple_oacc_kernels_set_clauses (p, clauses);
+
+  return p;
+}
+
+
 /* Build a GIMPLE_OACC_PARALLEL statement.
 
    BODY is sequence of statements which are executed in parallel.
@@ -1672,6 +1689,7 @@ gimple_copy (gimple stmt)
 	  gimple_try_set_cleanup (copy, new_seq);
 	  break;
 
+	case GIMPLE_OACC_KERNELS:
 	case GIMPLE_OACC_PARALLEL:
           gcc_unreachable ();
 
diff --git gcc/gimple.def gcc/gimple.def
index ce800bd..c9756b7 100644
--- gcc/gimple.def
+++ gcc/gimple.def
@@ -209,10 +209,28 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE)
    ordering is exposed by the range check in gimple_omp_subcode.  */
 
 
+/* GIMPLE_OACC_KERNELS <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
+   #pragma acc kernels [CLAUSES]
+   BODY is the sequence of statements inside the kernels construct.
+   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+   CHILD_FN is set when outlining the body of the kernels region.
+   All the statements in BODY are moved into this newly created
+   function when converting OMP constructs into low-GIMPLE.
+   DATA_ARG is a vec of 3 local variables in the parent function
+   containing data to be mapped to CHILD_FN.  This is used to
+   implement the MAP clauses.  */
+DEFGSCODE(GIMPLE_OACC_KERNELS, "gimple_oacc_kernels", GSS_OMP_PARALLEL_LAYOUT)
+
 /* GIMPLE_OACC_PARALLEL <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
-
    #pragma acc parallel [CLAUSES]
-   BODY */
+   BODY is the sequence of statements inside the parallel construct.
+   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+   CHILD_FN is set when outlining the body of the parallel region.
+   All the statements in BODY are moved into this newly created
+   function when converting OMP constructs into low-GIMPLE.
+   DATA_ARG is a vec of 3 local variables in the parent function
+   containing data to be mapped to CHILD_FN.  This is used to
+   implement the MAP clauses.  */
 DEFGSCODE(GIMPLE_OACC_PARALLEL, "gimple_oacc_parallel", GSS_OMP_PARALLEL_LAYOUT)
 
 /* Tuples used for lowering of OMP_ATOMIC.  Although the form of the OMP_ATOMIC
diff --git gcc/gimple.h gcc/gimple.h
index b4ee9fa..514af32 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -549,8 +549,8 @@ struct GTY((tag("GSS_OMP_FOR")))
 };
 
 
-/* GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL, GIMPLE_OMP_TARGET,
-   GIMPLE_OMP_TASK */
+/* GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL,
+   GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */
 struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
   gimple_statement_omp_parallel_layout : public gimple_statement_omp
 {
@@ -569,6 +569,14 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
   tree data_arg;
 };
 
+/* GIMPLE_OACC_KERNELS */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+  gimple_statement_oacc_kernels : public gimple_statement_omp_parallel_layout
+{
+    /* No extra fields; adds invariant:
+         stmt->code == GIMPLE_OACC_KERNELS.  */
+};
+
 /* GIMPLE_OACC_PARALLEL */
 struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
   gimple_statement_oacc_parallel : public gimple_statement_omp_parallel_layout
@@ -894,6 +902,14 @@ is_a_helper <gimple_statement_omp_for>::test (gimple gs)
 template <>
 template <>
 inline bool
+is_a_helper <gimple_statement_oacc_kernels>::test (gimple gs)
+{
+  return gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
 is_a_helper <gimple_statement_oacc_parallel>::test (gimple gs)
 {
   return gs->code == GIMPLE_OACC_PARALLEL;
@@ -1094,6 +1110,14 @@ is_a_helper <const gimple_statement_omp_for>::test (const_gimple gs)
 template <>
 template <>
 inline bool
+is_a_helper <const gimple_statement_oacc_kernels>::test (const_gimple gs)
+{
+  return gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
 is_a_helper <const gimple_statement_oacc_parallel>::test (const_gimple gs)
 {
   return gs->code == GIMPLE_OACC_PARALLEL;
@@ -1225,6 +1249,7 @@ gimple gimple_build_debug_bind_stat (tree, tree, gimple MEM_STAT_DECL);
 gimple gimple_build_debug_source_bind_stat (tree, tree, gimple MEM_STAT_DECL);
 #define gimple_build_debug_source_bind(var,val,stmt)			\
   gimple_build_debug_source_bind_stat ((var), (val), (stmt) MEM_STAT_INFO)
+gimple gimple_build_oacc_kernels (gimple_seq, tree);
 gimple gimple_build_oacc_parallel (gimple_seq, tree);
 gimple gimple_build_omp_critical (gimple_seq, tree);
 gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq);
@@ -1462,6 +1487,7 @@ gimple_has_substatements (gimple g)
     case GIMPLE_EH_FILTER:
     case GIMPLE_EH_ELSE:
     case GIMPLE_TRY:
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_FOR:
     case GIMPLE_OMP_MASTER:
@@ -4266,6 +4292,101 @@ gimple_omp_set_body (gimple gs, gimple_seq body)
 }
 
 
+/* Return the clauses associated with OACC_KERNELS statement GS.  */
+
+static inline tree
+gimple_oacc_kernels_clauses (const_gimple gs)
+{
+  const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <const gimple_statement_oacc_kernels> (gs);
+  return oacc_kernels_stmt->clauses;
+}
+
+/* Return a pointer to the clauses associated with OACC_KERNELS statement GS.  */
+
+static inline tree *
+gimple_oacc_kernels_clauses_ptr (gimple gs)
+{
+  gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <gimple_statement_oacc_kernels> (gs);
+  return &oacc_kernels_stmt->clauses;
+}
+
+/* Set CLAUSES to be the list of clauses associated with OACC_KERNELS statement
+   GS.  */
+
+static inline void
+gimple_oacc_kernels_set_clauses (gimple gs, tree clauses)
+{
+  gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <gimple_statement_oacc_kernels> (gs);
+  oacc_kernels_stmt->clauses = clauses;
+}
+
+/* Return the child function used to hold the body of OACC_KERNELS statement
+   GS.  */
+
+static inline tree
+gimple_oacc_kernels_child_fn (const_gimple gs)
+{
+  const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <const gimple_statement_oacc_kernels> (gs);
+  return oacc_kernels_stmt->child_fn;
+}
+
+/* Return a pointer to the child function used to hold the body of OACC_KERNELS
+   statement GS.  */
+
+static inline tree *
+gimple_oacc_kernels_child_fn_ptr (gimple gs)
+{
+  gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <gimple_statement_oacc_kernels> (gs);
+  return &oacc_kernels_stmt->child_fn;
+}
+
+/* Set CHILD_FN to be the child function for OACC_KERNELS statement GS.  */
+
+static inline void
+gimple_oacc_kernels_set_child_fn (gimple gs, tree child_fn)
+{
+  gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <gimple_statement_oacc_kernels> (gs);
+  oacc_kernels_stmt->child_fn = child_fn;
+}
+
+/* Return the artificial argument used to send variables and values
+   from the parent to the children threads in OACC_KERNELS statement GS.  */
+
+static inline tree
+gimple_oacc_kernels_data_arg (const_gimple gs)
+{
+  const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <const gimple_statement_oacc_kernels> (gs);
+  return oacc_kernels_stmt->data_arg;
+}
+
+/* Return a pointer to the data argument for OACC_KERNELS statement GS.  */
+
+static inline tree *
+gimple_oacc_kernels_data_arg_ptr (gimple gs)
+{
+  gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <gimple_statement_oacc_kernels> (gs);
+  return &oacc_kernels_stmt->data_arg;
+}
+
+/* Set DATA_ARG to be the data argument for OACC_KERNELS statement GS.  */
+
+static inline void
+gimple_oacc_kernels_set_data_arg (gimple gs, tree data_arg)
+{
+  gimple_statement_oacc_kernels *oacc_kernels_stmt =
+    as_a <gimple_statement_oacc_kernels> (gs);
+  oacc_kernels_stmt->data_arg = data_arg;
+}
+
+
 /* Return the clauses associated with OACC_PARALLEL statement GS.  */
 
 static inline tree
@@ -4330,7 +4451,8 @@ gimple_oacc_parallel_set_child_fn (gimple gs, tree child_fn)
   oacc_parallel_stmt->child_fn = child_fn;
 }
 
-/* Return the data argument for OACC_PARALLEL statement GS.  */
+/* Return the artificial argument used to send variables and values
+   from the parent to the children threads in OACC_PARALLEL statement GS.  */
 
 static inline tree
 gimple_oacc_parallel_data_arg (const_gimple gs)
@@ -5640,6 +5762,7 @@ gimple_return_set_retval (gimple gs, tree retval)
 /* Returns true when the gimple statement STMT is any of the OpenMP types.  */
 
 #define CASE_GIMPLE_OMP				\
+    case GIMPLE_OACC_KERNELS:			\
     case GIMPLE_OACC_PARALLEL:			\
     case GIMPLE_OMP_PARALLEL:			\
     case GIMPLE_OMP_TASK:			\
@@ -5683,6 +5806,7 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
   gcc_assert (is_gimple_omp (stmt));
   switch (gimple_code (stmt))
     {
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       return true;
     case GIMPLE_OMP_TARGET:
diff --git gcc/gimplify.c gcc/gimplify.c
index 6dbabfa..f3c34f9 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7007,6 +7007,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
       ort = (enum omp_region_type) (ORT_TARGET
 				    | ORT_TARGET_MAP_FORCE);
       break;
+    case OACC_KERNELS:
     case OACC_PARALLEL:
       ort = (enum omp_region_type) (ORT_TARGET
 				    | ORT_TARGET_OFFLOAD
@@ -7070,6 +7071,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
       stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
 				      OACC_DATA_CLAUSES (expr));
       break;
+    case OACC_KERNELS:
+      stmt = gimple_build_oacc_kernels (body, OACC_KERNELS_CLAUSES (expr));
+      break;
     case OACC_PARALLEL:
       stmt = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr));
       break;
@@ -8036,7 +8040,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = GS_ALL_DONE;
 	  break;
 
-	case OACC_KERNELS:
 	case OACC_HOST_DATA:
 	case OACC_DECLARE:
 	case OACC_UPDATE:
@@ -8066,6 +8069,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  break;
 
 	case OACC_DATA:
+	case OACC_KERNELS:
 	case OACC_PARALLEL:
 	case OMP_SECTIONS:
 	case OMP_SINGLE:
diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def
index eaf3228..2d5c22c 100644
--- gcc/oacc-builtins.def
+++ gcc/oacc-builtins.def
@@ -27,9 +27,11 @@ along with GCC; see the file COPYING3.  If not see
 
    See builtins.def for details.  */
 
-DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
-		   BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
 		   BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
 		   BT_FN_VOID, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_KERNELS, "GOACC_kernels",
+		   BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
+		   BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
diff --git gcc/omp-low.c gcc/omp-low.c
index eec862e..2f13fb4 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1502,6 +1502,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
   bool offloaded;
   switch (gimple_code (ctx->stmt))
     {
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       offloaded = true;
       break;
@@ -2085,13 +2086,28 @@ find_combined_for (gimple_stmt_iterator *gsi_p,
   return NULL;
 }
 
-/* Scan an OpenACC parallel directive.  */
+/* Scan an OpenACC offload directive.  */
 
 static void
-scan_oacc_parallel (gimple stmt, omp_context *outer_ctx)
+scan_oacc_offload (gimple stmt, omp_context *outer_ctx)
 {
   omp_context *ctx;
   tree name;
+  void (*gimple_omp_set_child_fn) (gimple, tree);
+  tree (*gimple_omp_clauses) (const_gimple);
+  switch (gimple_code (stmt))
+    {
+    case GIMPLE_OACC_KERNELS:
+      gimple_omp_set_child_fn = gimple_oacc_kernels_set_child_fn;
+      gimple_omp_clauses = gimple_oacc_kernels_clauses;
+      break;
+    case GIMPLE_OACC_PARALLEL:
+      gimple_omp_set_child_fn = gimple_oacc_parallel_set_child_fn;
+      gimple_omp_clauses = gimple_oacc_parallel_clauses;
+      break;
+    default:
+      gcc_unreachable ();
+    }
 
   gcc_assert (taskreg_nesting_level == 0);
   gcc_assert (target_nesting_level == 0);
@@ -2107,9 +2123,10 @@ scan_oacc_parallel (gimple stmt, omp_context *outer_ctx)
   DECL_NAMELESS (name) = 1;
   TYPE_NAME (ctx->record_type) = name;
   create_omp_child_function (ctx, false);
-  gimple_oacc_parallel_set_child_fn (stmt, ctx->cb.dst_fn);
 
-  scan_sharing_clauses (gimple_oacc_parallel_clauses (stmt), ctx);
+  gimple_omp_set_child_fn (stmt, ctx->cb.dst_fn);
+
+  scan_sharing_clauses (gimple_omp_clauses (stmt), ctx);
   scan_omp (gimple_omp_body_ptr (stmt), ctx);
 
   if (TYPE_FIELDS (ctx->record_type) == NULL)
@@ -2841,8 +2858,9 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 
   switch (gimple_code (stmt))
     {
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
-      scan_oacc_parallel (stmt, ctx);
+      scan_oacc_offload (stmt, ctx);
       break;
 
     case GIMPLE_OMP_PARALLEL:
@@ -4860,10 +4878,10 @@ expand_omp_build_assign (gimple_stmt_iterator *gsi_p, tree to, tree from)
     }
 }
 
-/* Expand the OpenACC parallel directive starting at REGION.  */
+/* Expand the OpenACC offload directive starting at REGION.  */
 
 static void
-expand_oacc_parallel (struct omp_region *region)
+expand_oacc_offload (struct omp_region *region)
 {
   basic_block entry_bb, exit_bb, new_bb;
   struct function *child_cfun;
@@ -4871,9 +4889,24 @@ expand_oacc_parallel (struct omp_region *region)
   gimple_stmt_iterator gsi;
   gimple entry_stmt, stmt;
   edge e;
+  tree (*gimple_omp_child_fn) (const_gimple);
+  tree (*gimple_omp_data_arg) (const_gimple);
+  switch (region->type)
+    {
+    case GIMPLE_OACC_KERNELS:
+      gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
+      gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+      break;
+    case GIMPLE_OACC_PARALLEL:
+      gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
+      gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
+      break;
+    default:
+      gcc_unreachable ();
+    }
 
   entry_stmt = last_stmt (region->entry);
-  child_fn = gimple_oacc_parallel_child_fn (entry_stmt);
+  child_fn = gimple_omp_child_fn (entry_stmt);
   child_cfun = DECL_STRUCT_FUNCTION (child_fn);
 
   /* Supported by expand_omp_taskreg, but not here.  */
@@ -4901,14 +4934,13 @@ expand_oacc_parallel (struct omp_region *region)
 	 a function call that has been inlined, the original PARM_DECL
 	 .OMP_DATA_I may have been converted into a different local
 	 variable.  In which case, we need to keep the assignment.  */
-      if (gimple_oacc_parallel_data_arg (entry_stmt))
+      if (gimple_omp_data_arg (entry_stmt))
 	{
 	  basic_block entry_succ_bb = single_succ (entry_bb);
 	  gimple_stmt_iterator gsi;
 	  tree arg;
 	  gimple parcopy_stmt = NULL;
-	  tree sender
-	    = TREE_VEC_ELT (gimple_oacc_parallel_data_arg (entry_stmt), 0);
+	  tree sender = TREE_VEC_ELT (gimple_omp_data_arg (entry_stmt), 0);
 
 	  for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
 	    {
@@ -4964,7 +4996,8 @@ expand_oacc_parallel (struct omp_region *region)
 	 so that it can be moved to the child function.  */
       gsi = gsi_last_bb (entry_bb);
       stmt = gsi_stmt (gsi);
-      gcc_assert (stmt && (gimple_code (stmt) == GIMPLE_OACC_PARALLEL));
+      gcc_assert (stmt && (gimple_code (stmt) == GIMPLE_OACC_KERNELS
+			   || gimple_code (stmt) == GIMPLE_OACC_PARALLEL));
       gsi_remove (&gsi, true);
       e = split_block (entry_bb, stmt);
       entry_bb = e->dest;
@@ -5037,10 +5070,22 @@ expand_oacc_parallel (struct omp_region *region)
   tree t1, t2, t3, t4, device, c, clauses;
   enum built_in_function start_ix;
   location_t clause_loc;
+  tree (*gimple_omp_clauses) (const_gimple);
+  switch (region->type)
+    {
+    case GIMPLE_OACC_KERNELS:
+      gimple_omp_clauses = gimple_oacc_kernels_clauses;
+      start_ix = BUILT_IN_GOACC_KERNELS;
+      break;
+    case GIMPLE_OACC_PARALLEL:
+      gimple_omp_clauses = gimple_oacc_parallel_clauses;
+      start_ix = BUILT_IN_GOACC_PARALLEL;
+      break;
+    default:
+      gcc_unreachable ();
+    }
 
-  clauses = gimple_oacc_parallel_clauses (entry_stmt);
-
-  start_ix = BUILT_IN_GOACC_PARALLEL;
+  clauses = gimple_omp_clauses (entry_stmt);
 
   /* By default, the value of DEVICE is -1 (let runtime library choose).  */
   device = build_int_cst (integer_type_node, -1);
@@ -5059,7 +5104,7 @@ expand_oacc_parallel (struct omp_region *region)
   device = fold_convert_loc (clause_loc, integer_type_node, device);
 
   gsi = gsi_last_bb (new_bb);
-  t = gimple_oacc_parallel_data_arg (entry_stmt);
+  t = gimple_omp_data_arg (entry_stmt);
   if (t == NULL)
     {
       t1 = size_zero_node;
@@ -8606,8 +8651,9 @@ expand_omp (struct omp_region *region)
 
       switch (region->type)
 	{
+	case GIMPLE_OACC_KERNELS:
 	case GIMPLE_OACC_PARALLEL:
-	  expand_oacc_parallel (region);
+	  expand_oacc_offload (region);
 	  break;
 
 	case GIMPLE_OMP_PARALLEL:
@@ -8851,11 +8897,11 @@ make_pass_expand_omp (gcc::context *ctxt)
 
 /* Routines to lower OpenMP directives into OMP-GIMPLE.  */
 
-/* Lower the OpenACC parallel directive in the current statement
+/* Lower the OpenACC offload directive in the current statement
    in GSI_P.  CTX holds context information for the directive.  */
 
 static void
-lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 {
   tree clauses;
   tree child_fn, t, c;
@@ -8864,8 +8910,23 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   gimple_seq par_body, olist, ilist, new_body;
   location_t loc = gimple_location (stmt);
   unsigned int map_cnt = 0;
+  tree (*gimple_omp_clauses) (const_gimple);
+  void (*gimple_omp_set_data_arg) (gimple, tree);
+  switch (gimple_code (stmt))
+    {
+    case GIMPLE_OACC_KERNELS:
+      gimple_omp_clauses = gimple_oacc_kernels_clauses;
+      gimple_omp_set_data_arg = gimple_oacc_kernels_set_data_arg;
+      break;
+    case GIMPLE_OACC_PARALLEL:
+      gimple_omp_clauses = gimple_oacc_parallel_clauses;
+      gimple_omp_set_data_arg = gimple_oacc_parallel_set_data_arg;
+      break;
+    default:
+      gcc_unreachable ();
+    }
 
-  clauses = gimple_oacc_parallel_clauses (stmt);
+  clauses = gimple_omp_clauses (stmt);
   par_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
   par_body = gimple_bind_body (par_bind);
   child_fn = ctx->cb.dst_fn;
@@ -8950,7 +9011,7 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
       TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
       TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1;
-      gimple_oacc_parallel_set_data_arg (stmt, t);
+      gimple_omp_set_data_arg (stmt, t);
 
       vec<constructor_elt, va_gc> *vsize;
       vec<constructor_elt, va_gc> *vkind;
@@ -10820,11 +10881,12 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     case GIMPLE_BIND:
       lower_omp (gimple_bind_body_ptr (stmt), ctx);
       break;
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       ctx = maybe_lookup_ctx (stmt);
       gcc_assert (ctx);
       gcc_assert (!ctx->cancellable);
-      lower_oacc_parallel (gsi_p, ctx);
+      lower_oacc_offload (gsi_p, ctx);
       break;
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
@@ -11053,6 +11115,9 @@ static bool
 diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
     	       gimple branch_ctx, gimple label_ctx)
 {
+  gcc_assert (!branch_ctx || is_gimple_omp (branch_ctx));
+  gcc_assert (!label_ctx || is_gimple_omp (label_ctx));
+
   if (label_ctx == branch_ctx)
     return false;
 
@@ -11070,8 +11135,8 @@ diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
     }
   if (flag_openacc)
     {
-      if ((branch_ctx && gimple_code (branch_ctx) == GIMPLE_OACC_PARALLEL)
-	  || (label_ctx && gimple_code (label_ctx) == GIMPLE_OACC_PARALLEL))
+      if ((branch_ctx && is_gimple_omp_oacc_specifically (branch_ctx))
+	  || (label_ctx && is_gimple_omp_oacc_specifically (label_ctx)))
 	{
 	  gcc_assert (kind == NULL);
 	  kind = "OpenACC";
@@ -11149,6 +11214,7 @@ diagnose_sb_1 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
     {
     WALK_SUBSTMTS;
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
@@ -11208,6 +11274,7 @@ diagnose_sb_2 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
     {
     WALK_SUBSTMTS;
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
@@ -11304,6 +11371,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
 
   switch (code)
     {
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
diff --git gcc/tree-inline.c gcc/tree-inline.c
index 61c1cc8..8b22b86 100644
--- gcc/tree-inline.c
+++ gcc/tree-inline.c
@@ -1316,6 +1316,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
 	  copy = gimple_build_wce (s1);
 	  break;
 
+	case GIMPLE_OACC_KERNELS:
 	case GIMPLE_OACC_PARALLEL:
           gcc_unreachable ();
 
@@ -3940,6 +3941,7 @@ estimate_num_insns (gimple stmt, eni_weights *weights)
               + estimate_num_insns_seq (gimple_omp_body (stmt), weights)
               + estimate_num_insns_seq (gimple_omp_for_pre_body (stmt), weights));
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
diff --git gcc/tree-nested.c gcc/tree-nested.c
index afa7abb..397f851 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1248,6 +1248,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	}
       break;
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       gcc_unreachable ();
 
@@ -1712,6 +1713,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 
   switch (gimple_code (stmt))
     {
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       gcc_unreachable ();
 
@@ -2075,6 +2077,7 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	break;
       }
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       gcc_unreachable ();
 
@@ -2138,6 +2141,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	}
       break;
 
+    case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
       gcc_unreachable ();
 
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index b90b09b..3ea5901 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2014-02-28  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* libgomp.map (GOACC_2.0): Add GOACC_kernels.
+	* libgomp_g.h (GOACC_kernels): New prototype.
+	* oacc-parallel.c (GOACC_kernels): New function.
+
 2014-02-21  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* testsuite/libgomp.oacc-c/data-1.c: New file.
diff --git libgomp/libgomp.map libgomp/libgomp.map
index cb52e45..e9f8b55 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -235,5 +235,6 @@ GOACC_2.0 {
   global:
 	GOACC_data_end;
 	GOACC_data_start;
+	GOACC_kernels;
 	GOACC_parallel;
 };
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index b9083a5..9681566 100644
--- libgomp/libgomp_g.h
+++ libgomp/libgomp_g.h
@@ -216,10 +216,12 @@ extern void GOMP_teams (unsigned int, unsigned int);
 
 /* oacc-parallel.c */
 
-extern void GOACC_parallel (int, void (*) (void *), const void *,
-			    size_t, void **, size_t *, unsigned short *);
 extern void GOACC_data_start (int, const void *,
 			      size_t, void **, size_t *, unsigned short *);
 extern void GOACC_data_end (void);
+extern void GOACC_kernels (int, void (*) (void *), const void *,
+			   size_t, void **, size_t *, unsigned short *);
+extern void GOACC_parallel (int, void (*) (void *), const void *,
+			    size_t, void **, size_t *, unsigned short *);
 
 #endif /* LIBGOMP_G_H */
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index 3ac7e39..cb883a8 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -23,7 +23,7 @@
    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
    <http://www.gnu.org/licenses/>.  */
 
-/* This file handles the OpenACC data and parallel constructs.  */
+/* This file handles OpenACC constructs.  */
 
 #include "libgomp.h"
 #include "libgomp_g.h"
@@ -81,3 +81,13 @@ GOACC_data_end (void)
 {
   GOMP_target_end_data ();
 }
+
+
+void
+GOACC_kernels (int device, void (*fn) (void *), const void *openmp_target,
+	       size_t mapnum, void **hostaddrs, size_t *sizes,
+	       unsigned short *kinds)
+{
+  /* TODO.  */
+  GOACC_parallel (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds);
+}
