diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index bece7c1..06e8583 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,4 +1,15 @@
 2014-12-10  Thomas Schwinge  <thomas@codesourcery.com>
+	    Bernd Schmidt  <bernds@codesourcery.com>
+
+	* gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge
+	into GIMPLE_OMP_TARGET.  Update all users.
+
+2014-12-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* cgraphbuild.c (pass_build_cgraph_edges::execute): Remove
+	handling of GIMPLE_OACC_PARALLEL.
+	* gimple-pretty-print.c (dump_gimple_omp_target): Dump a bit more
+	data, pretty-printing.
 
 	* omp-low.c (build_omp_regions_1, make_gimple_omp_edges)
 	<GIMPLE_OMP_TARGET>: Handle
diff --git gcc/cgraphbuild.c gcc/cgraphbuild.c
index 9b078bc..c72ceab 100644
--- gcc/cgraphbuild.c
+++ gcc/cgraphbuild.c
@@ -368,21 +368,14 @@ pass_build_cgraph_edges::execute (function *fun)
 					    bb->count, freq);
 	    }
 	  node->record_stmt_references (stmt);
-	  if (gimple_code (stmt) == GIMPLE_OACC_PARALLEL
-	      && gimple_oacc_parallel_child_fn (stmt))
-	    {
-	      tree fn = gimple_oacc_parallel_child_fn (stmt);
-	      node->create_reference (cgraph_node::get_create (fn),
-				      IPA_REF_ADDR, stmt);
-	    }
-	  else if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
-		   && gimple_omp_parallel_child_fn (stmt))
+	  if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
+	      && gimple_omp_parallel_child_fn (stmt))
 	    {
 	      tree fn = gimple_omp_parallel_child_fn (stmt);
 	      node->create_reference (cgraph_node::get_create (fn),
 				      IPA_REF_ADDR, stmt);
 	    }
-	  else if (gimple_code (stmt) == GIMPLE_OMP_TASK)
+	  if (gimple_code (stmt) == GIMPLE_OMP_TASK)
 	    {
 	      tree fn = gimple_omp_task_child_fn (stmt);
 	      if (fn)
diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi
index 4c59748..696c10e 100644
--- gcc/doc/gimple.texi
+++ gcc/doc/gimple.texi
@@ -439,8 +439,6 @@ 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
 @item @code{GIMPLE_OMP_CONTINUE}	@tab x			@tab x
@@ -1008,8 +1006,6 @@ 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}::
 * @code{GIMPLE_OMP_CONTINUE}::
@@ -1655,17 +1651,6 @@ Build a @code{GIMPLE_NOP} statement.
 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}
-
-
 @node @code{GIMPLE_OMP_ATOMIC_LOAD}
 @subsection @code{GIMPLE_OMP_ATOMIC_LOAD}
 @cindex @code{GIMPLE_OMP_ATOMIC_LOAD}
diff --git gcc/gimple-low.c gcc/gimple-low.c
index 60a7792e..3507d3c 100644
--- gcc/gimple-low.c
+++ gcc/gimple-low.c
@@ -368,8 +368,6 @@ 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:
     case GIMPLE_OMP_TARGET:
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index 72dfac6..38d39f7 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1338,15 +1338,21 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
     case GF_OMP_TARGET_KIND_UPDATE:
       kind = " update";
       break;
+    case GF_OMP_TARGET_KIND_OACC_KERNELS:
+      kind = " oacc_kernels";
+      break;
+    case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+      kind = " oacc_parallel";
+      break;
     case GF_OMP_TARGET_KIND_OACC_DATA:
       kind = " oacc_data";
       break;
-    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
-      kind = " oacc_enter_exit_data";
-      break;
     case GF_OMP_TARGET_KIND_OACC_UPDATE:
       kind = " oacc_update";
       break;
+    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+      kind = " oacc_enter_exit_data";
+      break;
     default:
       gcc_unreachable ();
     }
@@ -1355,7 +1361,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
       dump_gimple_fmt (buffer, spc, flags, "%G%s <%+BODY <%S>%nCLAUSES <", gs,
 		       kind, gimple_omp_body (gs));
       dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
-      dump_gimple_fmt (buffer, spc, flags, " >");
+      dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>",
+		       gimple_omp_target_child_fn (gs),
+		       gimple_omp_target_data_arg (gs));
     }
   else
     {
@@ -1367,16 +1375,28 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
 	  pp_string (buffer, " [child fn: ");
 	  dump_generic_node (buffer, gimple_omp_target_child_fn (gs),
 			     spc, flags, false);
-	  pp_right_bracket (buffer);
+	  pp_string (buffer, " (");
+	  if (gimple_omp_target_data_arg (gs))
+	    dump_generic_node (buffer, gimple_omp_target_data_arg (gs),
+			       spc, flags, false);
+	  else
+	    pp_string (buffer, "???");
+	  pp_string (buffer, ")]");
 	}
-      if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+      gimple_seq body = gimple_omp_body (gs);
+      if (body && gimple_code (gimple_seq_first_stmt (body)) != GIMPLE_BIND)
 	{
 	  newline_and_indent (buffer, spc + 2);
-	  pp_character (buffer, '{');
+	  pp_left_brace (buffer);
 	  pp_newline (buffer);
-	  dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+	  dump_gimple_seq (buffer, body, spc + 4, flags);
 	  newline_and_indent (buffer, spc + 2);
-	  pp_character (buffer, '}');
+	  pp_right_brace (buffer);
+	}
+      else if (body)
+	{
+	  pp_newline (buffer);
+	  dump_gimple_seq (buffer, body, spc + 2, flags);
 	}
     }
 }
@@ -1878,81 +1898,6 @@ dump_gimple_phi (pretty_printer *buffer, gimple phi, int spc, bool comment,
 }
 
 
-/* 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_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_omp_clauses (gs), spc, flags);
-      dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>",
-                       gimple_omp_child_fn (gs), gimple_omp_data_arg (gs));
-    }
-  else
-    {
-      gimple_seq body;
-      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_omp_child_fn (gs),
-			     spc, flags, false);
-	  pp_string (buffer, " (");
-	  if (gimple_omp_data_arg (gs))
-	    dump_generic_node (buffer, gimple_omp_data_arg (gs),
-			       spc, flags, false);
-	  else
-	    pp_string (buffer, "???");
-	  pp_string (buffer, ")]");
-	}
-      body = gimple_omp_body (gs);
-      if (body && gimple_code (gimple_seq_first_stmt (body)) != GIMPLE_BIND)
-	{
-	  newline_and_indent (buffer, spc + 2);
-	  pp_left_brace (buffer);
-	  pp_newline (buffer);
-	  dump_gimple_seq (buffer, body, spc + 4, flags);
-	  newline_and_indent (buffer, spc + 2);
-	  pp_right_brace (buffer);
-	}
-      else if (body)
-	{
-	  pp_newline (buffer);
-	  dump_gimple_seq (buffer, body, spc + 2, flags);
-	}
-    }
-}
-
-
 /* Dump a GIMPLE_OMP_PARALLEL tuple on the pretty_printer BUFFER, SPC spaces
    of indent.  FLAGS specifies details to show in the dump (see TDF_* in
    dumpfile.h).  */
@@ -2237,11 +2182,6 @@ 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_offload (buffer, gs, spc, flags);
-      break;
-
     case GIMPLE_OMP_PARALLEL:
       dump_gimple_omp_parallel (buffer, gs, spc, flags);
       break;
diff --git gcc/gimple-walk.c gcc/gimple-walk.c
index cc74d34..bfa3532 100644
--- gcc/gimple-walk.c
+++ gcc/gimple-walk.c
@@ -304,36 +304,6 @@ 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);
-      if (ret)
-	return ret;
-      ret = walk_tree (gimple_oacc_parallel_child_fn_ptr (stmt), callback_op,
-		       wi, pset);
-      if (ret)
-	return ret;
-      ret = walk_tree (gimple_oacc_parallel_data_arg_ptr (stmt), callback_op,
-		       wi, pset);
-      if (ret)
-	return ret;
-      break;
-
     case GIMPLE_OMP_CONTINUE:
       ret = walk_tree (gimple_omp_continue_control_def_ptr (stmt),
 	  	       callback_op, wi, pset);
@@ -629,8 +599,6 @@ 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:
     case GIMPLE_OMP_TASKGROUP:
diff --git gcc/gimple.c gcc/gimple.c
index e6de836..32615e8 100644
--- gcc/gimple.c
+++ gcc/gimple.c
@@ -811,40 +811,6 @@ 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.
-   CLAUSES are the OpenACC parallel construct's clauses.  */
-
-gimple
-gimple_build_oacc_parallel (gimple_seq body, tree clauses)
-{
-  gimple p = gimple_alloc (GIMPLE_OACC_PARALLEL, 0);
-  if (body)
-    gimple_omp_set_body (p, body);
-  gimple_oacc_parallel_set_clauses (p, clauses);
-
-  return p;
-}
-
-
 /* Build a GIMPLE_OMP_CRITICAL statement.
 
    BODY is the sequence of statements for which only one thread can execute.
@@ -1077,7 +1043,7 @@ gimple_build_omp_single (gimple_seq body, tree clauses)
 /* Build a GIMPLE_OMP_TARGET statement.
 
    BODY is the sequence of statements that will be executed.
-   KIND is the kind of target region.
+   KIND is the kind of the region.
    CLAUSES are any of the construct's clauses.  */
 
 gimple
@@ -1719,10 +1685,6 @@ gimple_copy (gimple stmt)
 	  gimple_try_set_cleanup (copy, new_seq);
 	  break;
 
-	case GIMPLE_OACC_KERNELS:
-	case GIMPLE_OACC_PARALLEL:
-          gcc_unreachable ();
-
 	case GIMPLE_OMP_FOR:
 	  gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
 	  new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt));
diff --git gcc/gimple.def gcc/gimple.def
index e2e912c..269c2d7 100644
--- gcc/gimple.def
+++ gcc/gimple.def
@@ -205,33 +205,8 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE)
 
 /* IMPORTANT.
 
-   Do not rearrange any of the GIMPLE_OACC_* and GIMPLE_OMP_* codes.  This
-   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 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)
+   Do not rearrange any of the GIMPLE_OMP_* codes.  This ordering is
+   exposed by the range check in gimple_omp_subcode.  */
 
 /* Tuples used for lowering of OMP_ATOMIC.  Although the form of the OMP_ATOMIC
    expression is very simple (just in form mem op= expr), various implicit
@@ -381,12 +356,12 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE)
 DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT)
 
 /* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
-   #pragma acc data
+   #pragma acc {kernels,parallel,data}
    #pragma omp target {,data,update}
-   BODY is the sequence of statements inside the target construct
+   BODY is the sequence of statements inside the construct
    (NULL for target update).
    CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
-   CHILD_FN is set when outlining the body of the target region.
+   CHILD_FN is set when outlining the body of the offloaded 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
diff --git gcc/gimple.h gcc/gimple.h
index a91bd4e..60c2469 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -108,9 +108,11 @@ enum gf_mask {
     GF_OMP_TARGET_KIND_REGION	= 0,
     GF_OMP_TARGET_KIND_DATA	= 1,
     GF_OMP_TARGET_KIND_UPDATE	= 2,
-    GF_OMP_TARGET_KIND_OACC_DATA = 3,
-    GF_OMP_TARGET_KIND_OACC_UPDATE = 4,
-    GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 5,
+    GF_OMP_TARGET_KIND_OACC_PARALLEL = 3,
+    GF_OMP_TARGET_KIND_OACC_KERNELS = 4,
+    GF_OMP_TARGET_KIND_OACC_DATA = 5,
+    GF_OMP_TARGET_KIND_OACC_UPDATE = 6,
+    GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 7,
 
     /* True on an GIMPLE_OMP_RETURN statement if the return does not require
        a thread synchronization via some sort of barrier.  The exact barrier
@@ -560,8 +562,8 @@ struct GTY((tag("GSS_OMP_FOR")))
 };
 
 
-/* GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL,
-   GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */
+/* 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
 {
@@ -580,22 +582,6 @@ 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
-{
-    /* No extra fields; adds invariant:
-         stmt->code == GIMPLE_OACC_PARALLEL.  */
-};
-
 /* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */
 struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
   gimple_statement_omp_taskreg : public gimple_statement_omp_parallel_layout
@@ -913,22 +899,6 @@ 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;
-}
-
-template <>
-template <>
-inline bool
 is_a_helper <gimple_statement_omp_taskreg *>::test (gimple gs)
 {
   return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
@@ -1121,22 +1091,6 @@ 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;
-}
-
-template <>
-template <>
-inline bool
 is_a_helper <const gimple_statement_omp_taskreg *>::test (const_gimple gs)
 {
   return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
@@ -1260,8 +1214,6 @@ 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);
 gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree);
@@ -1500,8 +1452,6 @@ 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:
     case GIMPLE_OMP_TASKGROUP:
@@ -4362,197 +4312,6 @@ 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
-gimple_oacc_parallel_clauses (const_gimple gs)
-{
-  const gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <const gimple_statement_oacc_parallel *> (gs);
-  return oacc_parallel_stmt->clauses;
-}
-
-/* Return a pointer to the clauses associated with OACC_PARALLEL statement
-   GS.  */
-
-static inline tree *
-gimple_oacc_parallel_clauses_ptr (gimple gs)
-{
-  gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <gimple_statement_oacc_parallel *> (gs);
-  return &oacc_parallel_stmt->clauses;
-}
-
-/* Set CLAUSES to be the list of clauses associated with OACC_PARALLEL
-   statement GS.  */
-
-static inline void
-gimple_oacc_parallel_set_clauses (gimple gs, tree clauses)
-{
-  gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <gimple_statement_oacc_parallel *> (gs);
-  oacc_parallel_stmt->clauses = clauses;
-}
-
-/* Return the child function used to hold the body of OACC_PARALLEL statement
-   GS.  */
-
-static inline tree
-gimple_oacc_parallel_child_fn (const_gimple gs)
-{
-  const gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <const gimple_statement_oacc_parallel *> (gs);
-  return oacc_parallel_stmt->child_fn;
-}
-
-/* Return a pointer to the child function used to hold the body of
-   OACC_PARALLEL statement GS.  */
-
-static inline tree *
-gimple_oacc_parallel_child_fn_ptr (gimple gs)
-{
-  gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <gimple_statement_oacc_parallel *> (gs);
-  return &oacc_parallel_stmt->child_fn;
-}
-
-/* Set CHILD_FN to be the child function for OACC_PARALLEL statement GS.  */
-
-static inline void
-gimple_oacc_parallel_set_child_fn (gimple gs, tree child_fn)
-{
-  gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <gimple_statement_oacc_parallel *> (gs);
-  oacc_parallel_stmt->child_fn = child_fn;
-}
-
-/* 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)
-{
-  const gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <const gimple_statement_oacc_parallel *> (gs);
-  return oacc_parallel_stmt->data_arg;
-}
-
-/* Return a pointer to the data argument for OACC_PARALLEL statement GS.  */
-
-static inline tree *
-gimple_oacc_parallel_data_arg_ptr (gimple gs)
-{
-  gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <gimple_statement_oacc_parallel *> (gs);
-  return &oacc_parallel_stmt->data_arg;
-}
-
-/* Set DATA_ARG to be the data argument for OACC_PARALLEL statement GS.  */
-
-static inline void
-gimple_oacc_parallel_set_data_arg (gimple gs, tree data_arg)
-{
-  gimple_statement_oacc_parallel *oacc_parallel_stmt =
-    as_a <gimple_statement_oacc_parallel *> (gs);
-  oacc_parallel_stmt->data_arg = data_arg;
-}
-
-
 /* Return the name associated with OMP_CRITICAL statement GS.  */
 
 static inline tree
@@ -5374,7 +5133,7 @@ gimple_omp_target_set_clauses (gimple gs, tree clauses)
 }
 
 
-/* Return the kind of OMP target statemement.  */
+/* Return the kind of the OMP_TARGET G.  */
 
 static inline int
 gimple_omp_target_kind (const_gimple g)
@@ -5384,7 +5143,7 @@ gimple_omp_target_kind (const_gimple g)
 }
 
 
-/* Set the OMP target kind.  */
+/* Set the kind of the OMP_TARGET G.  */
 
 static inline void
 gimple_omp_target_set_kind (gimple g, int kind)
@@ -5854,8 +5613,6 @@ gimple_return_set_retbnd (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:			\
     case GIMPLE_OMP_FOR:			\
@@ -5898,9 +5655,6 @@ 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_FOR:
       switch (gimple_omp_for_kind (stmt))
 	{
@@ -5908,10 +5662,12 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
 	  return true;
 	default:
 	  return false;
-	}      
+	}
     case GIMPLE_OMP_TARGET:
       switch (gimple_omp_target_kind (stmt))
 	{
+	case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+	case GF_OMP_TARGET_KIND_OACC_KERNELS:
 	case GF_OMP_TARGET_KIND_OACC_DATA:
 	case GF_OMP_TARGET_KIND_OACC_UPDATE:
 	case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
@@ -5933,13 +5689,12 @@ is_gimple_omp_offloaded (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:
       switch (gimple_omp_target_kind (stmt))
 	{
 	case GF_OMP_TARGET_KIND_REGION:
+	case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+	case GF_OMP_TARGET_KIND_OACC_KERNELS:
 	  return true;
 	default:
 	  return false;
diff --git gcc/gimplify.c gcc/gimplify.c
index ad48d51..eb9d930 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7315,10 +7315,12 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 				      OACC_DATA_CLAUSES (expr));
       break;
     case OACC_KERNELS:
-      stmt = gimple_build_oacc_kernels (body, OACC_KERNELS_CLAUSES (expr));
+      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
+				      OMP_CLAUSES (expr));
       break;
     case OACC_PARALLEL:
-      stmt = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr));
+      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
+				      OMP_CLAUSES (expr));
       break;
     case OMP_SECTIONS:
       stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
diff --git gcc/omp-low.c gcc/omp-low.c
index 6fed38f..39e2f22 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -272,10 +272,12 @@ oacc_max_threads (omp_context *ctx)
      Scan for the innermost vector_length clause.  */
   for (omp_context *oc = ctx; oc; oc = oc->outer)
     {
-      if (gimple_code (oc->stmt) != GIMPLE_OACC_PARALLEL)
+      if (gimple_code (oc->stmt) != GIMPLE_OMP_TARGET
+	  || (gimple_omp_target_kind (oc->stmt)
+	      != GF_OMP_TARGET_KIND_OACC_PARALLEL))
 	continue;
 
-      clauses = gimple_oacc_parallel_clauses (oc->stmt);
+      clauses = gimple_omp_target_clauses (oc->stmt);
 
       vector_length = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH);
       if (vector_length)
@@ -2643,28 +2645,7 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
 {
   omp_context *ctx;
   tree name;
-  bool offloaded;
-  void (*gimple_omp_set_child_fn) (gimple, tree);
-  tree (*gimple_omp_clauses) (const_gimple);
-
-  offloaded = is_gimple_omp_offloaded (stmt);
-  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;
-    case GIMPLE_OMP_TARGET:
-      gimple_omp_set_child_fn = gimple_omp_target_set_child_fn;
-      gimple_omp_clauses = gimple_omp_target_clauses;
-      break;
-    default:
-      gcc_unreachable ();
-    }
+  bool offloaded = is_gimple_omp_offloaded (stmt);
 
   if (is_gimple_omp_oacc_specifically (stmt))
     {
@@ -2689,10 +2670,10 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
 					     0, 0);
 
       create_omp_child_function (ctx, false);
-      gimple_omp_set_child_fn (stmt, ctx->cb.dst_fn);
+      gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
     }
 
-  scan_sharing_clauses (gimple_omp_clauses (stmt), ctx);
+  scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx);
   scan_omp (gimple_omp_body_ptr (stmt), ctx);
 
   if (TYPE_FIELDS (ctx->record_type) == NULL)
@@ -2944,8 +2925,6 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 		      "of work-sharing, critical, ordered, master or explicit "
 		      "task region");
 	    return false;
-	  case GIMPLE_OACC_KERNELS:
-	  case GIMPLE_OACC_PARALLEL:
 	  case GIMPLE_OMP_PARALLEL:
 	    return true;
 	  default:
@@ -3203,8 +3182,6 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       scan_omp (gimple_omp_body_ptr (stmt), ctx);
       break;
 
-    case GIMPLE_OACC_KERNELS:
-    case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_TARGET:
       scan_omp_target (stmt, ctx);
       break;
@@ -8780,42 +8757,24 @@ expand_omp_target (struct omp_region *region)
   gimple entry_stmt, stmt;
   edge e;
   bool offloaded, data_region;
-  tree (*gimple_omp_child_fn) (const_gimple);
-  tree (*gimple_omp_data_arg) (const_gimple);
 
   entry_stmt = last_stmt (region->entry);
   new_bb = region->entry;
 
   offloaded = is_gimple_omp_offloaded (entry_stmt);
-  data_region = false;
-  switch (region->type)
+  switch (gimple_omp_target_kind (entry_stmt))
     {
-    case GIMPLE_OACC_KERNELS:
-      gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
-      gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+    case GF_OMP_TARGET_KIND_REGION:
+    case GF_OMP_TARGET_KIND_UPDATE:
+    case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+    case GF_OMP_TARGET_KIND_OACC_KERNELS:
+    case GF_OMP_TARGET_KIND_OACC_UPDATE:
+    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+      data_region = false;
       break;
-    case GIMPLE_OACC_PARALLEL:
-      gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
-      gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
-      break;
-    case GIMPLE_OMP_TARGET:
-      switch (gimple_omp_target_kind (entry_stmt))
-	{
-	case GF_OMP_TARGET_KIND_DATA:
-	case GF_OMP_TARGET_KIND_OACC_DATA:
-	  data_region = true;
-	  break;
-	case GF_OMP_TARGET_KIND_REGION:
-	case GF_OMP_TARGET_KIND_UPDATE:
-	case GF_OMP_TARGET_KIND_OACC_UPDATE:
-	case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
-	  break;
-	default:
-	  gcc_unreachable ();
-	}
-
-      gimple_omp_child_fn = gimple_omp_target_child_fn;
-      gimple_omp_data_arg = gimple_omp_target_data_arg;
+    case GF_OMP_TARGET_KIND_DATA:
+    case GF_OMP_TARGET_KIND_OACC_DATA:
+      data_region = true;
       break;
     default:
       gcc_unreachable ();
@@ -8825,7 +8784,7 @@ expand_omp_target (struct omp_region *region)
   child_cfun = NULL;
   if (offloaded)
     {
-      child_fn = gimple_omp_child_fn (entry_stmt);
+      child_fn = gimple_omp_target_child_fn (entry_stmt);
       child_cfun = DECL_STRUCT_FUNCTION (child_fn);
     }
 
@@ -8854,13 +8813,14 @@ expand_omp_target (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_omp_data_arg (entry_stmt))
+      tree data_arg = gimple_omp_target_data_arg (entry_stmt);
+      if (data_arg)
 	{
 	  basic_block entry_succ_bb = single_succ (entry_bb);
 	  gimple_stmt_iterator gsi;
 	  tree arg;
 	  gimple tgtcopy_stmt = NULL;
-	  tree sender = TREE_VEC_ELT (gimple_omp_data_arg (entry_stmt), 0);
+	  tree sender = TREE_VEC_ELT (data_arg, 0);
 
 	  for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
 	    {
@@ -8994,49 +8954,38 @@ expand_omp_target (struct omp_region *region)
   tree t1, t2, t3, t4, device, cond, c, clauses;
   enum built_in_function start_ix;
   location_t clause_loc;
-  tree (*gimple_omp_clauses) (const_gimple);
 
-  switch (region->type)
+  switch (gimple_omp_target_kind (entry_stmt))
     {
-    case GIMPLE_OACC_KERNELS:
-      gimple_omp_clauses = gimple_oacc_kernels_clauses;
-      start_ix = BUILT_IN_GOACC_KERNELS;
+    case GF_OMP_TARGET_KIND_REGION:
+      start_ix = BUILT_IN_GOMP_TARGET;
       break;
-    case GIMPLE_OACC_PARALLEL:
-      gimple_omp_clauses = gimple_oacc_parallel_clauses;
+    case GF_OMP_TARGET_KIND_DATA:
+      start_ix = BUILT_IN_GOMP_TARGET_DATA;
+      break;
+    case GF_OMP_TARGET_KIND_UPDATE:
+      start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+      break;
+    case GF_OMP_TARGET_KIND_OACC_PARALLEL:
       start_ix = BUILT_IN_GOACC_PARALLEL;
       break;
-    case GIMPLE_OMP_TARGET:
-      gimple_omp_clauses = gimple_omp_target_clauses;
-      switch (gimple_omp_target_kind (entry_stmt))
-	{
-	case GF_OMP_TARGET_KIND_REGION:
-	  start_ix = BUILT_IN_GOMP_TARGET;
-	  break;
-	case GF_OMP_TARGET_KIND_DATA:
-	  start_ix = BUILT_IN_GOMP_TARGET_DATA;
-	  break;
-	case GF_OMP_TARGET_KIND_UPDATE:
-	  start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
-	  break;
-	case GF_OMP_TARGET_KIND_OACC_DATA:
-	  start_ix = BUILT_IN_GOACC_DATA_START;
-	  break;
-	case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
-	  start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
-	  break;
-	case GF_OMP_TARGET_KIND_OACC_UPDATE:
-	  start_ix = BUILT_IN_GOACC_UPDATE;
-	  break;
-	default:
-	  gcc_unreachable ();
-	}
+    case GF_OMP_TARGET_KIND_OACC_KERNELS:
+      start_ix = BUILT_IN_GOACC_KERNELS;
+      break;
+    case GF_OMP_TARGET_KIND_OACC_DATA:
+      start_ix = BUILT_IN_GOACC_DATA_START;
+      break;
+    case GF_OMP_TARGET_KIND_OACC_UPDATE:
+      start_ix = BUILT_IN_GOACC_UPDATE;
+      break;
+    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+      start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
       break;
     default:
       gcc_unreachable ();
     }
 
-  clauses = gimple_omp_clauses (entry_stmt);
+  clauses = gimple_omp_target_clauses (entry_stmt);
 
   /* By default, the value of DEVICE is -1 (let runtime library choose)
      and there is no conditional.  */
@@ -9119,7 +9068,7 @@ expand_omp_target (struct omp_region *region)
     }
 
   gsi = gsi_last_bb (new_bb);
-  t = gimple_omp_data_arg (entry_stmt);
+  t = gimple_omp_target_data_arg (entry_stmt);
   if (t == NULL)
     {
       t1 = size_zero_node;
@@ -9331,8 +9280,6 @@ expand_omp (struct omp_region *region)
 	  expand_omp_atomic (region);
 	  break;
 
-	case GIMPLE_OACC_KERNELS:
-	case GIMPLE_OACC_PARALLEL:
 	case GIMPLE_OMP_TARGET:
 	  expand_omp_target (region);
 	  break;
@@ -9679,22 +9626,18 @@ oacc_initialize_reduction_data (tree clauses, tree nthreads,
   tree c, t, oc;
   gimple stmt;
   omp_context *octx;
-  tree (*gimple_omp_clauses) (const_gimple);
-  void (*gimple_omp_set_clauses) (gimple, tree);
 
   /* Find the innermost PARALLEL openmp context.  FIXME: OpenACC kernels
      may require extra care unless they are converted to openmp for loops.  */
-
-  if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL)
+  if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+      && (gimple_omp_target_kind (ctx->stmt)
+	  == GF_OMP_TARGET_KIND_OACC_PARALLEL))
     octx = ctx;
   else
     octx = ctx->outer;
 
-  gimple_omp_clauses = gimple_oacc_parallel_clauses;
-  gimple_omp_set_clauses = gimple_oacc_parallel_set_clauses;
-
   /* Extract the clauses.  */
-  oc = gimple_omp_clauses (octx->stmt);
+  oc = gimple_omp_target_clauses (octx->stmt);
 
   /* Find the last outer clause.  */
   for (; oc && OMP_CLAUSE_CHAIN (oc); oc = OMP_CLAUSE_CHAIN (oc))
@@ -9744,7 +9687,7 @@ oacc_initialize_reduction_data (tree clauses, tree nthreads,
       if (oc)
 	OMP_CLAUSE_CHAIN (oc) = t;
       else
-	gimple_omp_set_clauses (octx->stmt, t);
+	gimple_omp_target_set_clauses (octx->stmt, t);
       OMP_CLAUSE_SIZE (t) = size;
       oc = t;
     }
@@ -11195,45 +11138,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   location_t loc = gimple_location (stmt);
   bool offloaded, data_region;
   unsigned int map_cnt = 0;
-  tree (*gimple_omp_clauses) (const_gimple);
-  void (*gimple_omp_set_data_arg) (gimple, tree);
 
   offloaded = is_gimple_omp_offloaded (stmt);
-  data_region = false;
-  switch (gimple_code (stmt))
+  switch (gimple_omp_target_kind (stmt))
     {
-    case GIMPLE_OACC_KERNELS:
-      gimple_omp_clauses = gimple_oacc_kernels_clauses;
-      gimple_omp_set_data_arg = gimple_oacc_kernels_set_data_arg;
+    case GF_OMP_TARGET_KIND_REGION:
+    case GF_OMP_TARGET_KIND_UPDATE:
+    case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+    case GF_OMP_TARGET_KIND_OACC_KERNELS:
+    case GF_OMP_TARGET_KIND_OACC_UPDATE:
+    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+      data_region = false;
       break;
-    case GIMPLE_OACC_PARALLEL:
-      gimple_omp_clauses = gimple_oacc_parallel_clauses;
-      gimple_omp_set_data_arg = gimple_oacc_parallel_set_data_arg;
-      break;
-    case GIMPLE_OMP_TARGET:
-      switch (gimple_omp_target_kind (stmt))
-	{
-	case GF_OMP_TARGET_KIND_DATA:
-	case GF_OMP_TARGET_KIND_OACC_DATA:
-	  data_region = true;
-	  break;
-	case GF_OMP_TARGET_KIND_REGION:
-	case GF_OMP_TARGET_KIND_UPDATE:
-	case GF_OMP_TARGET_KIND_OACC_UPDATE:
-	case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
-	  break;
-	default:
-	  gcc_unreachable ();
-	}
-
-      gimple_omp_clauses = gimple_omp_target_clauses;
-      gimple_omp_set_data_arg = gimple_omp_target_set_data_arg;
+    case GF_OMP_TARGET_KIND_DATA:
+    case GF_OMP_TARGET_KIND_OACC_DATA:
+      data_region = true;
       break;
     default:
       gcc_unreachable ();
     }
 
-  clauses = gimple_omp_clauses (stmt);
+  clauses = gimple_omp_target_clauses (stmt);
 
   tgt_bind = NULL;
   tgt_body = NULL;
@@ -11250,15 +11175,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
   irlist = NULL;
   orlist = NULL;
-  switch (gimple_code (stmt))
-    {
-    case GIMPLE_OACC_KERNELS:
-    case GIMPLE_OACC_PARALLEL:
-      oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx);
-      break;
-    default:
-      break;
-    }
+  if (offloaded
+      && is_gimple_omp_oacc_specifically (stmt))
+    oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx);
 
   for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
     switch (OMP_CLAUSE_CODE (c))
@@ -11390,7 +11309,7 @@ lower_omp_target (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_omp_set_data_arg (stmt, t);
+      gimple_omp_target_set_data_arg (stmt, t);
 
       vec<constructor_elt, va_gc> *vsize;
       vec<constructor_elt, va_gc> *vkind;
@@ -11457,8 +11376,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 			    || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
 		if (maybe_lookup_oacc_reduction (var, ctx))
 		  {
-		    gcc_assert (gimple_code (stmt) == GIMPLE_OACC_KERNELS
-				|| gimple_code (stmt) == GIMPLE_OACC_PARALLEL);
+		    gcc_assert (offloaded
+				&& is_gimple_omp_oacc_specifically (stmt));
 		    gimplify_assign (x, var, &ilist);
 		  }
 		else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -11802,8 +11721,6 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 			lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
 	gimple_regimplify_operands (stmt, gsi_p);
       break;
-    case GIMPLE_OACC_KERNELS:
-    case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_TARGET:
       ctx = maybe_lookup_ctx (stmt);
       gcc_assert (ctx);
@@ -12095,8 +12012,6 @@ 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:
     case GIMPLE_OMP_SECTIONS:
@@ -12155,8 +12070,6 @@ 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:
     case GIMPLE_OMP_SECTIONS:
@@ -12252,8 +12165,6 @@ 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:
     case GIMPLE_OMP_FOR:
diff --git gcc/testsuite/gfortran.dg/goacc/private-1.f95 gcc/testsuite/gfortran.dg/goacc/private-1.f95
index 54c027d..23ce95a 100644
--- gcc/testsuite/gfortran.dg/goacc/private-1.f95
+++ gcc/testsuite/gfortran.dg/goacc/private-1.f95
@@ -31,7 +31,7 @@ program test
   end do
   !$acc end parallel
 end program test
-! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel" 3 "omplower" } }
 ! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
 ! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
diff --git gcc/tree-inline.c gcc/tree-inline.c
index 54b3514..89cb1eb 100644
--- gcc/tree-inline.c
+++ gcc/tree-inline.c
@@ -1395,10 +1395,6 @@ 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 ();
-
 	case GIMPLE_OMP_PARALLEL:
 	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
 	  copy = gimple_build_omp_parallel
@@ -4112,8 +4108,6 @@ 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:
     case GIMPLE_OMP_CRITICAL:
diff --git gcc/tree-nested.c gcc/tree-nested.c
index b5d6543..4da6297 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1325,10 +1325,6 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	}
       break;
 
-    case GIMPLE_OACC_KERNELS:
-    case GIMPLE_OACC_PARALLEL:
-      gcc_unreachable ();
-
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_suppress = info->suppress_expansion;
@@ -1898,10 +1894,6 @@ 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 ();
-
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_suppress = info->suppress_expansion;
@@ -2289,10 +2281,6 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	break;
       }
 
-    case GIMPLE_OACC_KERNELS:
-    case GIMPLE_OACC_PARALLEL:
-      gcc_unreachable ();
-
     case GIMPLE_OMP_TARGET:
       gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
@@ -2360,10 +2348,6 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	}
       break;
 
-    case GIMPLE_OACC_KERNELS:
-    case GIMPLE_OACC_PARALLEL:
-      gcc_unreachable ();
-
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_static_chain_added = info->static_chain_added;
