diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c
index f001a75..f7d2bd9 100644
--- gcc/c-family/c-omp.c
+++ gcc/c-family/c-omp.c
@@ -627,6 +627,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
   enum c_omp_clause_split s;
   int i;
 
+  gcc_assert (code != OACC_PARALLEL);
   for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
     cclauses[i] = NULL;
   /* Add implicit nowait clause on
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index 98f98d0..c329f8d 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1165,6 +1165,7 @@ static vec<pragma_ns_name> registered_pp_pragmas;
 
 struct omp_pragma_def { const char *name; unsigned int id; };
 static const struct omp_pragma_def oacc_pragmas[] = {
+  { "parallel", PRAGMA_OACC_PARALLEL },
 };
 static const struct omp_pragma_def omp_pragmas[] = {
   { "atomic", PRAGMA_OMP_ATOMIC },
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index 705bcb4..5c58e32 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -27,6 +27,7 @@ along with GCC; see the file COPYING3.  If not see
 typedef enum pragma_kind {
   PRAGMA_NONE = 0,
 
+  PRAGMA_OACC_PARALLEL,
   PRAGMA_OMP_ATOMIC,
   PRAGMA_OMP_BARRIER,
   PRAGMA_OMP_CANCEL,
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 8a1e988..297b6da7 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -4478,6 +4478,17 @@ c_parser_label (c_parser *parser)
      @throw expression ;
      @throw ;
 
+   OpenACC:
+
+   statement:
+     openacc-construct
+
+   openacc-construct:
+     parallel-construct
+
+   parallel-construct:
+     parallel-directive structured-block
+
    OpenMP:
 
    statement:
@@ -10754,7 +10765,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
   return clauses;
 }
 
-/* OpenMP 2.5:
+/* OpenACC 2.0, OpenMP 2.5:
    structured-block:
      statement
 
@@ -10770,6 +10781,32 @@ c_parser_omp_structured_block (c_parser *parser)
   return pop_stmt_list (stmt);
 }
 
+/* OpenACC 2.0:
+   # pragma acc parallel oacc-parallel-clause[optseq] new-line
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_PARALLEL_CLAUSE_MASK			\
+	PRAGMA_OMP_CLAUSE_NONE
+
+static tree
+c_parser_oacc_parallel (location_t loc, c_parser *parser)
+{
+  tree stmt, clauses, block;
+
+  clauses =  c_parser_omp_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
+				       "#pragma acc parallel");
+  gcc_assert (clauses == NULL);
+
+  block = c_begin_omp_parallel ();
+  add_stmt (c_parser_omp_structured_block (parser));
+
+  stmt = c_finish_oacc_parallel (loc, clauses, block);
+
+  return stmt;
+}
+
 /* OpenMP 2.5:
    # pragma omp atomic new-line
      expression-stmt
@@ -12948,6 +12985,9 @@ c_parser_omp_construct (c_parser *parser)
 
   switch (p_kind)
     {
+    case PRAGMA_OACC_PARALLEL:
+      stmt = c_parser_oacc_parallel (loc, parser);
+      break;
     case PRAGMA_OMP_ATOMIC:
       c_parser_omp_atomic (loc, parser);
       return;
diff --git gcc/c/c-tree.h gcc/c/c-tree.h
index 2565ccb..f524e31 100644
--- gcc/c/c-tree.h
+++ gcc/c/c-tree.h
@@ -635,6 +635,7 @@ extern tree c_finish_bc_stmt (location_t, tree *, bool);
 extern tree c_finish_goto_label (location_t, tree);
 extern tree c_finish_goto_ptr (location_t, tree);
 extern tree c_expr_to_decl (tree, bool *, bool *);
+extern tree c_finish_oacc_parallel (location_t, tree, tree);
 extern tree c_begin_omp_parallel (void);
 extern tree c_finish_omp_parallel (location_t, tree, tree);
 extern tree c_begin_omp_task (void);
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index 8f1d3a4..e7096e6 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -10644,6 +10644,25 @@ c_expr_to_decl (tree expr, bool *tc ATTRIBUTE_UNUSED, bool *se)
     return expr;
 }
 
+/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
+   statement.  LOC is the location of the OACC_PARALLEL.  */
+
+tree
+c_finish_oacc_parallel (location_t loc, tree clauses, tree block)
+{
+  tree stmt;
+
+  block = c_end_compound_stmt (loc, block, true);
+
+  stmt = make_node (OACC_PARALLEL);
+  TREE_TYPE (stmt) = void_type_node;
+  OACC_PARALLEL_CLAUSES (stmt) = clauses;
+  OACC_PARALLEL_BODY (stmt) = block;
+  SET_EXPR_LOCATION (stmt, loc);
+
+  return add_stmt (stmt);
+}
+
 /* Like c_begin_compound_stmt, except force the retention of the BLOCK.  */
 
 tree
diff --git gcc/cgraphbuild.c gcc/cgraphbuild.c
index 87e06e3..efad3d9 100644
--- gcc/cgraphbuild.c
+++ gcc/cgraphbuild.c
@@ -333,7 +333,15 @@ build_cgraph_edges (void)
 					     bb->count, freq);
 	    }
 	  ipa_record_stmt_references (node, stmt);
-	  if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
+	  if (gimple_code (stmt) == GIMPLE_OACC_PARALLEL
+	      && gimple_oacc_parallel_child_fn (stmt))
+	    {
+	      tree fn = gimple_oacc_parallel_child_fn (stmt);
+	      ipa_record_reference (node,
+				    cgraph_get_create_real_symbol_node (fn),
+				    IPA_REF_ADDR, stmt);
+	    }
+	  else if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
 	      && gimple_omp_parallel_child_fn (stmt))
 	    {
 	      tree fn = gimple_omp_parallel_child_fn (stmt);
@@ -341,7 +349,7 @@ build_cgraph_edges (void)
 				    cgraph_get_create_real_symbol_node (fn),
 				    IPA_REF_ADDR, stmt);
 	    }
-	  if (gimple_code (stmt) == GIMPLE_OMP_TASK)
+	  else if (gimple_code (stmt) == GIMPLE_OMP_TASK)
 	    {
 	      tree fn = gimple_omp_task_child_fn (stmt);
 	      if (fn)
diff --git gcc/doc/generic.texi gcc/doc/generic.texi
index 73dd123..812f5a9 100644
--- gcc/doc/generic.texi
+++ gcc/doc/generic.texi
@@ -2049,6 +2049,7 @@ edge.  Rethrowing the exception is represented using @code{RESX_EXPR}.
 
 @node OpenMP
 @subsection OpenMP
+@tindex OACC_PARALLEL
 @tindex OMP_PARALLEL
 @tindex OMP_FOR
 @tindex OMP_SECTIONS
@@ -2066,6 +2067,10 @@ All the statements starting with @code{OMP_} represent directives and
 clauses used by the OpenMP API @w{@uref{http://www.openmp.org/}}.
 
 @table @code
+@item OACC_PARALLEL
+
+Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}.
+
 @item OMP_PARALLEL
 
 Represents @code{#pragma omp parallel [clause1 @dots{} clauseN]}. It
diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi
index 7bd9fd5..0f1bbe6 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_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
@@ -905,6 +906,7 @@ Return a deep copy of statement @code{STMT}.
 * @code{GIMPLE_EH_FILTER}::
 * @code{GIMPLE_LABEL}::
 * @code{GIMPLE_NOP}::
+* @code{GIMPLE_OACC_PARALLEL}::
 * @code{GIMPLE_OMP_ATOMIC_LOAD}::
 * @code{GIMPLE_OMP_ATOMIC_STORE}::
 * @code{GIMPLE_OMP_CONTINUE}::
@@ -1554,6 +1556,12 @@ Build a @code{GIMPLE_NOP} statement.
 Returns @code{TRUE} if statement @code{G} is a @code{GIMPLE_NOP}.
 @end deftypefn
 
+
+@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 d527d86..74c9925 100644
--- gcc/gimple-low.c
+++ gcc/gimple-low.c
@@ -368,6 +368,7 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
       }
       break;
 
+    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 6842213..59cb5bb 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1823,6 +1823,60 @@ 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
+   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)
+{
+  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_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>",
+                       gimple_oacc_parallel_child_fn (gs),
+                       gimple_oacc_parallel_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, " [child fn: ");
+	  dump_generic_node (buffer, gimple_oacc_parallel_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),
+			       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).  */
@@ -2123,6 +2177,10 @@ 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_PARALLEL:
+      dump_gimple_oacc_parallel (buffer, gs, spc, flags);
+      break;
+
     case GIMPLE_OMP_PARALLEL:
       dump_gimple_omp_parallel (buffer, gs, spc, flags);
       break;
diff --git gcc/gimple.c gcc/gimple.c
index 20f6010..ea96d26 100644
--- gcc/gimple.c
+++ gcc/gimple.c
@@ -898,6 +898,23 @@ gimple_build_debug_source_bind_stat (tree var, tree value,
 }
 
 
+/* 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.
@@ -1571,6 +1588,21 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op,
 	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);
@@ -1866,6 +1898,7 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
 	return wi->callback_result;
 
       /* FALL THROUGH.  */
+    case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_CRITICAL:
     case GIMPLE_OMP_MASTER:
     case GIMPLE_OMP_TASKGROUP:
@@ -2306,6 +2339,9 @@ gimple_copy (gimple stmt)
 	  gimple_try_set_cleanup (copy, new_seq);
 	  break;
 
+	case GIMPLE_OACC_PARALLEL:
+          abort ();
+
 	case GIMPLE_OMP_FOR:
 	  new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt));
 	  gimple_omp_for_set_pre_body (copy, new_seq);
diff --git gcc/gimple.def gcc/gimple.def
index 07370ae..9ff9ab3 100644
--- gcc/gimple.def
+++ gcc/gimple.def
@@ -205,10 +205,16 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE)
 
 /* IMPORTANT.
 
-   Do not rearrange any of the GIMPLE_OMP_* codes.  This ordering is
-   exposed by the range check in gimple_omp_subcode().  */
+   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_PARALLEL <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
+
+   #pragma acc parallel [CLAUSES]
+   BODY */
+DEFGSCODE(GIMPLE_OACC_PARALLEL, "gimple_oacc_parallel", GSS_OMP_PARALLEL)
+
 /* 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
    conversions may cause the expression to become more complex, so that it does
diff --git gcc/gimple.h gcc/gimple.h
index b34424c..c9be1c9 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -786,6 +786,7 @@ gimple gimple_build_resx (int);
 gimple gimple_build_eh_dispatch (int);
 gimple gimple_build_switch_nlabels (unsigned, tree, tree);
 gimple gimple_build_switch (tree, tree, vec<tree> );
+gimple gimple_build_oacc_parallel (gimple_seq, tree);
 gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree);
 gimple gimple_build_omp_task (gimple_seq, tree, tree, tree, tree, tree, tree);
 gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq);
@@ -1256,6 +1257,7 @@ gimple_has_substatements (gimple g)
     case GIMPLE_EH_FILTER:
     case GIMPLE_EH_ELSE:
     case GIMPLE_TRY:
+    case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_FOR:
     case GIMPLE_OMP_MASTER:
     case GIMPLE_OMP_TASKGROUP:
@@ -4061,6 +4063,92 @@ gimple_omp_set_body (gimple gs, gimple_seq body)
 }
 
 
+/* Return the clauses associated with OACC_PARALLEL statement GS.  */
+
+static inline tree
+gimple_oacc_parallel_clauses (const_gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  return gs->gimple_omp_parallel.clauses;
+}
+
+/* Return a pointer to the clauses associated with OACC_PARALLEL statement
+   GS.  */
+
+static inline tree *
+gimple_oacc_parallel_clauses_ptr (gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  return &gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  gs->gimple_omp_parallel.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)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  return gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  return &gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  gs->gimple_omp_parallel.child_fn = child_fn;
+}
+
+/* Return the data argument for OACC_PARALLEL statement GS.  */
+
+static inline tree
+gimple_oacc_parallel_data_arg (const_gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  return gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  return &gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL);
+  gs->gimple_omp_parallel.data_arg = data_arg;
+}
+
+
 /* Return the name associated with OMP_CRITICAL statement GS.  */
 
 static inline tree
@@ -5269,6 +5357,7 @@ 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_PARALLEL:			\
     case GIMPLE_OMP_PARALLEL:			\
     case GIMPLE_OMP_TASK:			\
     case GIMPLE_OMP_FOR:			\
diff --git gcc/gimplify.c gcc/gimplify.c
index 30c2b45..0c45729 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -4641,6 +4641,7 @@ is_gimple_stmt (tree t)
     case CATCH_EXPR:
     case ASM_EXPR:
     case STATEMENT_LIST:
+    case OACC_PARALLEL:
     case OMP_PARALLEL:
     case OMP_FOR:
     case OMP_SIMD:
@@ -6745,6 +6746,37 @@ gimplify_adjust_omp_clauses (tree *list_p)
   delete_omp_context (ctx);
 }
 
+/* Gimplify the contents of an OACC_PARALLEL statement.  This involves
+   gimplification of the body, as well as scanning the body for used
+   variables.  We need to do this scan now, because variable-sized
+   decls will be decomposed during gimplification.  */
+
+static void
+gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p)
+{
+  tree expr = *expr_p;
+  gimple g;
+  gimple_seq body = NULL;
+  struct gimplify_ctx gctx;
+
+  gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p,
+			     ORT_TARGET);
+
+  push_gimplify_context (&gctx);
+
+  g = gimplify_and_return_first (OACC_PARALLEL_BODY (expr), &body);
+  if (gimple_code (g) == GIMPLE_BIND)
+    pop_gimplify_context (g);
+  else
+    pop_gimplify_context (NULL);
+
+  gimplify_adjust_omp_clauses (&OACC_PARALLEL_CLAUSES (expr));
+
+  g = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr));
+  gimplify_seq_add_stmt (pre_p, g);
+  *expr_p = NULL_TREE;
+}
+
 /* Gimplify the contents of an OMP_PARALLEL statement.  This involves
    gimplification of the body, as well as scanning the body for used
    variables.  We need to do this scan now, because variable-sized
@@ -8169,6 +8201,11 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = GS_ALL_DONE;
 	  break;
 
+	case OACC_PARALLEL:
+	  gimplify_oacc_parallel (expr_p, pre_p);
+	  ret = GS_ALL_DONE;
+	  break;
+
 	case OMP_PARALLEL:
 	  gimplify_omp_parallel (expr_p, pre_p);
 	  ret = GS_ALL_DONE;
@@ -8575,6 +8612,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 		  && code != LOOP_EXPR
 		  && code != SWITCH_EXPR
 		  && code != TRY_FINALLY_EXPR
+		  && code != OACC_PARALLEL
 		  && code != OMP_CRITICAL
 		  && code != OMP_FOR
 		  && code != OMP_MASTER
diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def
index fd630e0..a75e42d 100644
--- gcc/oacc-builtins.def
+++ gcc/oacc-builtins.def
@@ -26,3 +26,6 @@ along with GCC; see the file COPYING3.  If not see
      DEF_GOACC_BUILTIN (ENUM, NAME, TYPE, ATTRS)
 
    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)
diff --git gcc/omp-low.c gcc/omp-low.c
index 99811d0..84fe466 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -844,6 +844,8 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx)
      when we know the value is not accessible from an outer scope.  */
   if (shared_ctx)
     {
+      gcc_assert (gimple_code (shared_ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
       /* ??? Trivially accessible from anywhere.  But why would we even
 	 be passing an address in this case?  Should we simply assert
 	 this to be false, or should we have a cleanup pass that removes
@@ -985,6 +987,8 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
 static tree
 build_outer_var_ref (tree var, omp_context *ctx)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   tree x;
 
   if (is_global_var (maybe_lookup_decl_in_outer_ctx (var, ctx)))
@@ -1484,6 +1488,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
       switch (OMP_CLAUSE_CODE (c))
 	{
 	case OMP_CLAUSE_PRIVATE:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (OMP_CLAUSE_PRIVATE_OUTER_REF (c))
 	    goto do_private;
@@ -1492,6 +1497,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE_SHARED:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  /* Ignore shared directives in teams construct.  */
 	  if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
 	    break;
@@ -1518,6 +1524,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  goto do_private;
 
 	case OMP_CLAUSE_LASTPRIVATE:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  /* Let the corresponding firstprivate clause create
 	     the variable.  */
 	  if (OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c))
@@ -1527,6 +1534,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_FIRSTPRIVATE:
 	case OMP_CLAUSE_REDUCTION:
 	case OMP_CLAUSE_LINEAR:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  decl = OMP_CLAUSE_DECL (c);
 	do_private:
 	  if (is_variable_sized (decl))
@@ -1555,6 +1563,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE__LOOPTEMP_:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  gcc_assert (is_parallel_ctx (ctx));
 	  decl = OMP_CLAUSE_DECL (c);
 	  install_var_field (decl, false, 3, ctx);
@@ -1563,12 +1572,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 
 	case OMP_CLAUSE_COPYPRIVATE:
 	case OMP_CLAUSE_COPYIN:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  decl = OMP_CLAUSE_DECL (c);
 	  by_ref = use_pointer_for_field (decl, NULL);
 	  install_var_field (decl, by_ref, 3, ctx);
 	  break;
 
 	case OMP_CLAUSE_DEFAULT:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
 	  break;
 
@@ -1581,6 +1592,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_SCHEDULE:
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_DEPEND:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  if (ctx->outer)
 	    scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
 	  break;
@@ -1599,10 +1611,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && lookup_attribute ("omp declare target",
 				   DECL_ATTRIBUTES (decl)))
+	    {
+	      gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	    break;
+	    }
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
 	    {
+	      gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	      /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
 		 #pragma omp target data, there is nothing to map for
 		 those.  */
@@ -1632,8 +1648,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		    install_var_field (decl, true, 7, ctx);
 		  else
 		    install_var_field (decl, true, 3, ctx);
-		  if (gimple_omp_target_kind (ctx->stmt)
-		      == GF_OMP_TARGET_KIND_REGION)
+		  if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL
+		      || (gimple_omp_target_kind (ctx->stmt)
+			  == GF_OMP_TARGET_KIND_REGION))
 		    install_var_local (decl, ctx);
 		}
 	    }
@@ -1673,9 +1690,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_MERGEABLE:
 	case OMP_CLAUSE_PROC_BIND:
 	case OMP_CLAUSE_SAFELEN:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (is_global_var (decl)
 	      && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
@@ -1692,6 +1711,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
       switch (OMP_CLAUSE_CODE (c))
 	{
 	case OMP_CLAUSE_LASTPRIVATE:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  /* Let the corresponding firstprivate clause create
 	     the variable.  */
 	  if (OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c))
@@ -1704,6 +1724,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_FIRSTPRIVATE:
 	case OMP_CLAUSE_REDUCTION:
 	case OMP_CLAUSE_LINEAR:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (is_variable_sized (decl))
 	    install_var_local (decl, ctx);
@@ -1716,6 +1737,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE_SHARED:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  /* Ignore shared directives in teams construct.  */
 	  if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
 	    break;
@@ -1725,14 +1747,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE_MAP:
-	  if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+	  if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+	      && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
 	    break;
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (DECL_P (decl)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && lookup_attribute ("omp declare target",
 				   DECL_ATTRIBUTES (decl)))
+	    {
+	      gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	    break;
+	    }
 	  if (DECL_P (decl))
 	    {
 	      if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
@@ -1781,6 +1807,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE__LOOPTEMP_:
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
+	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
 	  break;
 
 	default:
@@ -1789,6 +1816,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
     }
 
   if (scan_array_reductions)
+    {
+      gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
     for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
 	  && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
@@ -1799,6 +1828,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
       else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
 	       && OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c))
 	scan_omp (&OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c), ctx);
+    }
 }
 
 /* Create a new name for omp child function.  Returns an identifier.  */
@@ -1830,6 +1860,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
   decl = build_decl (gimple_location (ctx->stmt),
 		     FUNCTION_DECL, name, type);
 
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+	      || !task_copy);
   if (!task_copy)
     ctx->cb.dst_fn = decl;
   else
@@ -1861,6 +1893,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
 	    break;
 	  }
     }
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+	      || !target_p);
   if (target_p)
     DECL_ATTRIBUTES (decl)
       = tree_cons (get_identifier ("omp declare target"),
@@ -1935,6 +1969,52 @@ find_combined_for (gimple_stmt_iterator *gsi_p,
   return NULL;
 }
 
+/* Scan an OpenACC parallel directive.  */
+
+static void
+scan_oacc_parallel (gimple stmt, omp_context *outer_ctx)
+{
+  omp_context *ctx;
+  tree name;
+
+  gcc_assert (taskreg_nesting_level == 0);
+  gcc_assert (target_nesting_level == 0);
+
+  ctx = new_omp_context (stmt, outer_ctx);
+  ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
+  ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
+  ctx->record_type = lang_hooks.types.make_type (RECORD_TYPE);
+  name = create_tmp_var_name (".omp_data_t");
+  name = build_decl (gimple_location (stmt),
+		     TYPE_DECL, name, ctx->record_type);
+  DECL_ARTIFICIAL (name) = 1;
+  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);
+  scan_omp (gimple_omp_body_ptr (stmt), ctx);
+
+  if (TYPE_FIELDS (ctx->record_type) == NULL)
+    ctx->record_type = ctx->receiver_decl = NULL;
+  else
+    {
+      TYPE_FIELDS (ctx->record_type)
+	= nreverse (TYPE_FIELDS (ctx->record_type));
+#ifdef ENABLE_CHECKING
+      tree field;
+      unsigned int align = DECL_ALIGN (TYPE_FIELDS (ctx->record_type));
+      for (field = TYPE_FIELDS (ctx->record_type);
+	   field;
+	   field = DECL_CHAIN (field))
+	gcc_assert (DECL_ALIGN (field) == align);
+#endif
+      layout_type (ctx->record_type);
+      fixup_child_record_type (ctx);
+    }
+}
+
 /* Scan an OpenMP parallel directive.  */
 
 static void
@@ -2225,6 +2305,38 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx)
 static bool
 check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 {
+  omp_context *ctx_;
+
+  /* TODO: While the OpenACC specification does allow for certain kinds of
+     nesting, we don't support that yet.  */
+  /* No nesting of STMT (which is an OpenACC or OpenMP one, or a GOMP builtin)
+     inside any OpenACC CTX.  */
+  for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+    switch (gimple_code (ctx_->stmt))
+      {
+      case GIMPLE_OACC_PARALLEL:
+	error_at (gimple_location (stmt),
+		  "may not be nested");
+	return false;
+      default:
+	break;
+      }
+  /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX.  */
+  switch (gimple_code (stmt))
+    {
+    case GIMPLE_OACC_PARALLEL:
+      for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+	if (is_gimple_omp (ctx_->stmt))
+	  {
+	    error_at (gimple_location (stmt),
+		      "may not be nested");
+	    return false;
+	  }
+      break;
+    default:
+      break;
+    }
+
   if (ctx != NULL)
     {
       if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
@@ -2584,6 +2696,10 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 
   switch (gimple_code (stmt))
     {
+    case GIMPLE_OACC_PARALLEL:
+      scan_oacc_parallel (stmt, ctx);
+      break;
+
     case GIMPLE_OMP_PARALLEL:
       taskreg_nesting_level++;
       scan_omp_parallel (gsi, ctx);
@@ -2910,6 +3026,8 @@ static bool
 lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, int &max_vf,
 			      tree &idx, tree &lane, tree &ivar, tree &lvar)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   if (max_vf == 0)
     {
       max_vf = omp_max_vf ();
@@ -2959,6 +3077,8 @@ static void
 lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
 			 omp_context *ctx, struct omp_for_data *fd)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   tree c, dtor, copyin_seq, x, ptr;
   bool copyin_by_ref = false;
   bool lastprivate_firstprivate = false;
@@ -3617,6 +3737,8 @@ static void
 lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
 			   omp_context *ctx)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   tree x, c, label = NULL, orig_clauses = clauses;
   bool par_clauses = false;
   tree simduid = NULL, lastlane = NULL;
@@ -3752,6 +3874,8 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
 static void
 lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   gimple_seq sub_seq = NULL;
   gimple stmt;
   tree x, c;
@@ -3853,6 +3977,8 @@ static void
 lower_copyprivate_clauses (tree clauses, gimple_seq *slist, gimple_seq *rlist,
 			    omp_context *ctx)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   tree c;
 
   for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -3903,6 +4029,8 @@ static void
 lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist,
     		    omp_context *ctx)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   tree c;
 
   for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -3994,6 +4122,8 @@ lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist,
 static void
 lower_send_shared_vars (gimple_seq *ilist, gimple_seq *olist, omp_context *ctx)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   tree var, ovar, nvar, f, x, record_type;
 
   if (ctx->record_type == NULL)
@@ -4542,6 +4672,234 @@ expand_omp_build_assign (gimple_stmt_iterator *gsi_p, tree to, tree from)
     }
 }
 
+/* Expand the OpenACC parallel directive starting at REGION.  */
+
+static void
+expand_oacc_parallel (struct omp_region *region)
+{
+  basic_block entry_bb, exit_bb, new_bb;
+  struct function *child_cfun;
+  tree child_fn, block, t;
+  gimple_stmt_iterator gsi;
+  gimple entry_stmt, stmt;
+  edge e;
+
+  entry_stmt = last_stmt (region->entry);
+  child_fn = gimple_oacc_parallel_child_fn (entry_stmt);
+  child_cfun = DECL_STRUCT_FUNCTION (child_fn);
+
+  /* Supported by expand_omp_taskreg, but not here.  */
+  gcc_assert (!child_cfun->cfg);
+  gcc_assert (!gimple_in_ssa_p (cfun));
+
+  entry_bb = region->entry;
+  exit_bb = region->exit;
+
+  /* Preserve indentation of expand_omp_target and expand_omp_taskreg.  */
+  if (1)
+    {
+      unsigned srcidx, dstidx, num;
+
+      /* If the parallel region needs data sent from the parent
+	 function, then the very first statement (except possible
+	 tree profile counter updates) of the parallel body
+	 is a copy assignment .OMP_DATA_I = &.OMP_DATA_O.  Since
+	 &.OMP_DATA_O is passed as an argument to the child function,
+	 we need to replace it with the argument as seen by the child
+	 function.
+
+	 In most cases, this will end up being the identity assignment
+	 .OMP_DATA_I = .OMP_DATA_I.  However, if the parallel body had
+	 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))
+	{
+	  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);
+
+	  for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
+	    {
+	      gcc_assert (!gsi_end_p (gsi));
+	      stmt = gsi_stmt (gsi);
+	      if (gimple_code (stmt) != GIMPLE_ASSIGN)
+		continue;
+
+	      if (gimple_num_ops (stmt) == 2)
+		{
+		  tree arg = gimple_assign_rhs1 (stmt);
+
+		  /* We're ignore the subcode because we're
+		     effectively doing a STRIP_NOPS.  */
+
+		  if (TREE_CODE (arg) == ADDR_EXPR
+		      && TREE_OPERAND (arg, 0) == sender)
+		    {
+		      parcopy_stmt = stmt;
+		      break;
+		    }
+		}
+	    }
+
+	  gcc_assert (parcopy_stmt != NULL);
+	  arg = DECL_ARGUMENTS (child_fn);
+
+	  gcc_assert (gimple_assign_lhs (parcopy_stmt) == arg);
+	  gsi_remove (&gsi, true);
+	}
+
+      /* Declare local variables needed in CHILD_CFUN.  */
+      block = DECL_INITIAL (child_fn);
+      BLOCK_VARS (block) = vec2chain (child_cfun->local_decls);
+      /* The gimplifier could record temporaries in the block
+	 rather than in containing function's local_decls chain,
+	 which would mean cgraph missed finalizing them.  Do it now.  */
+      for (t = BLOCK_VARS (block); t; t = DECL_CHAIN (t))
+	if (TREE_CODE (t) == VAR_DECL
+	    && TREE_STATIC (t)
+	    && !DECL_EXTERNAL (t))
+	  varpool_finalize_decl (t);
+      DECL_SAVED_TREE (child_fn) = NULL;
+      /* We'll create a CFG for child_fn, so no gimple body is needed.  */
+      gimple_set_body (child_fn, NULL);
+      TREE_USED (block) = 1;
+
+      /* Reset DECL_CONTEXT on function arguments.  */
+      for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t))
+	DECL_CONTEXT (t) = child_fn;
+
+      /* Split ENTRY_BB at GIMPLE_OACC_PARALLEL,
+	 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));
+      gsi_remove (&gsi, true);
+      e = split_block (entry_bb, stmt);
+      entry_bb = e->dest;
+      single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
+
+      /* Convert GIMPLE_OMP_RETURN into a RETURN_EXPR.  */
+      if (exit_bb)
+	{
+	  gsi = gsi_last_bb (exit_bb);
+	  gcc_assert (!gsi_end_p (gsi)
+		      && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
+	  stmt = gimple_build_return (NULL);
+	  gsi_insert_after (&gsi, stmt, GSI_SAME_STMT);
+	  gsi_remove (&gsi, true);
+	}
+
+      /* Move the region into CHILD_CFUN.  */
+
+      block = gimple_block (entry_stmt);
+
+      new_bb = move_sese_region_to_fn (child_cfun, entry_bb, exit_bb, block);
+      if (exit_bb)
+	single_succ_edge (new_bb)->flags = EDGE_FALLTHRU;
+      /* When the expansion process cannot guarantee an up-to-date
+         loop tree arrange for the child function to fixup loops.  */
+      if (loops_state_satisfies_p (LOOPS_NEED_FIXUP))
+	child_cfun->x_current_loops->state |= LOOPS_NEED_FIXUP;
+
+      /* Remove non-local VAR_DECLs from child_cfun->local_decls list.  */
+      num = vec_safe_length (child_cfun->local_decls);
+      for (srcidx = 0, dstidx = 0; srcidx < num; srcidx++)
+	{
+	  t = (*child_cfun->local_decls)[srcidx];
+	  if (DECL_CONTEXT (t) == cfun->decl)
+	    continue;
+	  if (srcidx != dstidx)
+	    (*child_cfun->local_decls)[dstidx] = t;
+	  dstidx++;
+	}
+      if (dstidx != num)
+	vec_safe_truncate (child_cfun->local_decls, dstidx);
+
+      /* Inform the callgraph about the new function.  */
+      DECL_STRUCT_FUNCTION (child_fn)->curr_properties = cfun->curr_properties;
+      cgraph_add_new_function (child_fn, true);
+
+      /* Fix the callgraph edges for child_cfun.  Those for cfun will be
+	 fixed in a following pass.  */
+      push_cfun (child_cfun);
+      rebuild_cgraph_edges ();
+
+      /* Some EH regions might become dead, see PR34608.  If
+	 pass_cleanup_cfg isn't the first pass to happen with the
+	 new child, these dead EH edges might cause problems.
+	 Clean them up now.  */
+      if (flag_exceptions)
+	{
+	  basic_block bb;
+	  bool changed = false;
+
+	  FOR_EACH_BB (bb)
+	    changed |= gimple_purge_dead_eh_edges (bb);
+	  if (changed)
+	    cleanup_tree_cfg ();
+	}
+      pop_cfun ();
+    }
+
+  /* Emit a library call to launch CHILD_FN.  */
+  tree t1, t2, t3, t4, device, c, clauses;
+  enum built_in_function start_ix;
+  location_t clause_loc;
+
+  clauses = gimple_oacc_parallel_clauses (entry_stmt);
+
+  start_ix = BUILT_IN_GOACC_PARALLEL;
+
+  /* By default, the value of DEVICE is -1 (let runtime library choose).  */
+  device = build_int_cst (integer_type_node, -1);
+
+  c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
+  gcc_assert (c == NULL);
+  if (c)
+    {
+      device = OMP_CLAUSE_DEVICE_ID (c);
+      clause_loc = OMP_CLAUSE_LOCATION (c);
+    }
+  else
+    clause_loc = gimple_location (entry_stmt);
+
+  /* Ensure 'device' is of the correct type.  */
+  device = fold_convert_loc (clause_loc, integer_type_node, device);
+
+  gsi = gsi_last_bb (new_bb);
+  t = gimple_oacc_parallel_data_arg (entry_stmt);
+  if (t == NULL)
+    {
+      t1 = size_zero_node;
+      t2 = build_zero_cst (ptr_type_node);
+      t3 = t2;
+      t4 = t2;
+    }
+  else
+    {
+      t1 = TYPE_MAX_VALUE (TYPE_DOMAIN (TREE_TYPE (TREE_VEC_ELT (t, 1))));
+      t1 = size_binop (PLUS_EXPR, t1, size_int (1));
+      t2 = build_fold_addr_expr (TREE_VEC_ELT (t, 0));
+      t3 = build_fold_addr_expr (TREE_VEC_ELT (t, 1));
+      t4 = build_fold_addr_expr (TREE_VEC_ELT (t, 2));
+    }
+
+  gimple g;
+  /* FIXME: This will be address of
+     extern char __OPENMP_TARGET__[] __attribute__((visibility ("hidden")))
+     symbol, as soon as the linker plugin is able to create it for us.  */
+  tree openmp_target = build_zero_cst (ptr_type_node);
+  tree fnaddr = build_fold_addr_expr (child_fn);
+  g = gimple_build_call (builtin_decl_explicit (start_ix),
+			 7, device, fnaddr, openmp_target, t1, t2, t3, t4);
+  gimple_set_location (g, gimple_location (entry_stmt));
+  gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+}
+
 /* Expand the OpenMP parallel or task directive starting at REGION.  */
 
 static void
@@ -8037,6 +8395,10 @@ expand_omp (struct omp_region *region)
 
       switch (region->type)
 	{
+	case GIMPLE_OACC_PARALLEL:
+	  expand_oacc_parallel (region);
+	  break;
+
 	case GIMPLE_OMP_PARALLEL:
 	case GIMPLE_OMP_TASK:
 	  expand_omp_taskreg (region);
@@ -8278,6 +8640,288 @@ make_pass_expand_omp (gcc::context *ctxt)
 
 /* Routines to lower OpenMP directives into OMP-GIMPLE.  */
 
+/* Lower the OpenACC parallel 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)
+{
+  tree clauses;
+  tree child_fn, t, c;
+  gimple stmt = gsi_stmt (*gsi_p);
+  gimple par_bind, bind;
+  gimple_seq par_body, olist, ilist, new_body;
+  struct gimplify_ctx gctx;
+  location_t loc = gimple_location (stmt);
+  unsigned int map_cnt = 0;
+
+  clauses = gimple_oacc_parallel_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;
+
+  push_gimplify_context (&gctx);
+
+  for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+    switch (OMP_CLAUSE_CODE (c))
+      {
+	tree var, x;
+
+      default:
+	break;
+      case OMP_CLAUSE_MAP:
+      case OMP_CLAUSE_TO:
+      case OMP_CLAUSE_FROM:
+	var = OMP_CLAUSE_DECL (c);
+	if (!DECL_P (var))
+	  {
+	    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+		|| !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+	      map_cnt++;
+	    continue;
+	  }
+
+	if (DECL_SIZE (var)
+	    && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+	  {
+	    tree var2 = DECL_VALUE_EXPR (var);
+	    gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
+	    var2 = TREE_OPERAND (var2, 0);
+	    gcc_assert (DECL_P (var2));
+	    var = var2;
+	  }
+
+	if (!maybe_lookup_field (var, ctx))
+	  continue;
+
+	/* Preserve indentation of lower_omp_target.  */
+	if (1)
+	  {
+	    x = build_receiver_ref (var, true, ctx);
+	    tree new_var = lookup_decl (var, ctx);
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+		&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+		&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+	      x = build_simple_mem_ref (x);
+	    SET_DECL_VALUE_EXPR (new_var, x);
+	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	  }
+	map_cnt++;
+      }
+
+  target_nesting_level++;
+  lower_omp (&par_body, ctx);
+  target_nesting_level--;
+
+  /* Declare all the variables created by mapping and the variables
+     declared in the scope of the body.  */
+  record_vars_into (ctx->block_vars, child_fn);
+  record_vars_into (gimple_bind_vars (par_bind), child_fn);
+
+  olist = NULL;
+  ilist = NULL;
+  if (ctx->record_type)
+    {
+      ctx->sender_decl
+	= create_tmp_var (ctx->record_type, ".omp_data_arr");
+      DECL_NAMELESS (ctx->sender_decl) = 1;
+      TREE_ADDRESSABLE (ctx->sender_decl) = 1;
+      t = make_tree_vec (3);
+      TREE_VEC_ELT (t, 0) = ctx->sender_decl;
+      TREE_VEC_ELT (t, 1)
+	= create_tmp_var (build_array_type_nelts (size_type_node, map_cnt),
+			  ".omp_data_sizes");
+      DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
+      TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
+      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
+      TREE_VEC_ELT (t, 2)
+	= create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
+						  map_cnt),
+			  ".omp_data_kinds");
+      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);
+
+      vec<constructor_elt, va_gc> *vsize;
+      vec<constructor_elt, va_gc> *vkind;
+      vec_alloc (vsize, map_cnt);
+      vec_alloc (vkind, map_cnt);
+      unsigned int map_idx = 0;
+
+      for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+	switch (OMP_CLAUSE_CODE (c))
+	  {
+	    tree ovar, nc;
+
+	  default:
+	    break;
+	  case OMP_CLAUSE_MAP:
+	  case OMP_CLAUSE_TO:
+	  case OMP_CLAUSE_FROM:
+	    nc = c;
+	    ovar = OMP_CLAUSE_DECL (c);
+	    if (!DECL_P (ovar))
+	      {
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+		  {
+		    gcc_checking_assert (OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (c))
+					 == get_base_address (ovar));
+		    nc = OMP_CLAUSE_CHAIN (c);
+		    ovar = OMP_CLAUSE_DECL (nc);
+		  }
+		else
+		  {
+		    tree x = build_sender_ref (ovar, ctx);
+		    tree v
+		      = build_fold_addr_expr_with_type (ovar, ptr_type_node);
+		    gimplify_assign (x, v, &ilist);
+		    nc = NULL_TREE;
+		  }
+	      }
+	    else
+	      {
+		if (DECL_SIZE (ovar)
+		    && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
+		  {
+		    tree ovar2 = DECL_VALUE_EXPR (ovar);
+		    gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
+		    ovar2 = TREE_OPERAND (ovar2, 0);
+		    gcc_assert (DECL_P (ovar2));
+		    ovar = ovar2;
+		  }
+		if (!maybe_lookup_field (ovar, ctx))
+		  continue;
+	      }
+
+	    if (nc)
+	      {
+		tree var = lookup_decl_in_outer_ctx (ovar, ctx);
+		tree x = build_sender_ref (ovar, ctx);
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+		    && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+		    && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+		  {
+		    tree avar
+		      = create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL);
+		    mark_addressable (avar);
+		    gimplify_assign (avar, build_fold_addr_expr (var), &ilist);
+		    avar = build_fold_addr_expr (avar);
+		    gimplify_assign (x, avar, &ilist);
+		  }
+		else if (is_gimple_reg (var))
+		  {
+		    tree avar = create_tmp_var (TREE_TYPE (var), NULL);
+		    mark_addressable (avar);
+		    if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
+			&& OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
+		      gimplify_assign (avar, var, &ilist);
+		    avar = build_fold_addr_expr (avar);
+		    gimplify_assign (x, avar, &ilist);
+		    if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
+			 || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
+			&& !TYPE_READONLY (TREE_TYPE (var)))
+		      {
+			x = build_sender_ref (ovar, ctx);
+			x = build_simple_mem_ref (x);
+			gimplify_assign (var, x, &olist);
+		      }
+		  }
+		else
+		  {
+		    var = build_fold_addr_expr (var);
+		    gimplify_assign (x, var, &ilist);
+		  }
+	      }
+	    tree s = OMP_CLAUSE_SIZE (c);
+	    if (s == NULL_TREE)
+	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
+	    s = fold_convert (size_type_node, s);
+	    tree purpose = size_int (map_idx++);
+	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
+	    if (TREE_CODE (s) != INTEGER_CST)
+	      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
+
+	    unsigned char tkind = 0;
+	    switch (OMP_CLAUSE_CODE (c))
+	      {
+	      case OMP_CLAUSE_MAP:
+		tkind = OMP_CLAUSE_MAP_KIND (c);
+		break;
+	      case OMP_CLAUSE_TO:
+		tkind = OMP_CLAUSE_MAP_TO;
+		break;
+	      case OMP_CLAUSE_FROM:
+		tkind = OMP_CLAUSE_MAP_FROM;
+		break;
+	      default:
+		gcc_unreachable ();
+	      }
+	    unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+	    if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
+	      talign = DECL_ALIGN_UNIT (ovar);
+	    talign = ceil_log2 (talign);
+	    tkind |= talign << 3;
+	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
+				    build_int_cst (unsigned_char_type_node,
+						   tkind));
+	    if (nc && nc != c)
+	      c = nc;
+	  }
+
+      gcc_assert (map_idx == map_cnt);
+
+      DECL_INITIAL (TREE_VEC_ELT (t, 1))
+	= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
+      DECL_INITIAL (TREE_VEC_ELT (t, 2))
+	= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind);
+      if (!TREE_STATIC (TREE_VEC_ELT (t, 1)))
+	{
+	  gimple_seq initlist = NULL;
+	  force_gimple_operand (build1 (DECL_EXPR, void_type_node,
+					TREE_VEC_ELT (t, 1)),
+				&initlist, true, NULL_TREE);
+	  gimple_seq_add_seq (&ilist, initlist);
+	}
+
+      tree clobber = build_constructor (ctx->record_type, NULL);
+      TREE_THIS_VOLATILE (clobber) = 1;
+      gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
+							clobber));
+    }
+
+  /* Once all the expansions are done, sequence all the different
+     fragments inside gimple_omp_body.  */
+
+  new_body = NULL;
+
+  if (ctx->record_type)
+    {
+      t = build_fold_addr_expr_loc (loc, ctx->sender_decl);
+      /* fixup_child_record_type might have changed receiver_decl's type.  */
+      t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t);
+      gimple_seq_add_stmt (&new_body,
+			   gimple_build_assign (ctx->receiver_decl, t));
+    }
+
+  gimple_seq_add_seq (&new_body, par_body);
+  gcc_assert (!ctx->cancellable);
+  new_body = maybe_catch_exception (new_body);
+  gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
+  gimple_omp_set_body (stmt, new_body);
+
+  bind = gimple_build_bind (NULL, NULL, gimple_bind_block (par_bind));
+  gsi_replace (gsi_p, bind, true);
+  gimple_bind_add_seq (bind, ilist);
+  gimple_bind_add_stmt (bind, stmt);
+  gimple_bind_add_seq (bind, olist);
+
+  pop_gimplify_context (NULL);
+}
+
 /* If ctx is a worksharing context inside of a cancellable parallel
    region and it isn't nowait, add lhs to its GIMPLE_OMP_RETURN
    and conditional branch to parallel's cancel_label to handle
@@ -8286,6 +8930,8 @@ make_pass_expand_omp (gcc::context *ctxt)
 static void
 maybe_add_implicit_barrier_cancel (omp_context *ctx, gimple_seq *body)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   gimple omp_return = gimple_seq_last_stmt (*body);
   gcc_assert (gimple_code (omp_return) == GIMPLE_OMP_RETURN);
   if (gimple_omp_return_nowait_p (omp_return))
@@ -9051,6 +9697,8 @@ task_copyfn_remap_type (struct omp_taskcopy_context *tcctx, tree orig_type)
 static void
 create_task_copyfn (gimple task_stmt, omp_context *ctx)
 {
+  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
   struct function *child_cfun;
   tree child_fn, t, c, src, dst, f, sf, arg, sarg, decl;
   tree record_type, srecord_type, bind, list;
@@ -9909,6 +10557,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_PARALLEL:
+      ctx = maybe_lookup_ctx (stmt);
+      gcc_assert (ctx);
+      gcc_assert (!ctx->cancellable);
+      lower_oacc_parallel (gsi_p, ctx);
+      break;
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       ctx = maybe_lookup_ctx (stmt);
@@ -10357,6 +11011,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region)
 
   switch (code)
     {
+    case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
     case GIMPLE_OMP_FOR:
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
new file mode 100644
index 0000000..875ec66
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -0,0 +1,121 @@
+/* TODO: Some of these should either be allowed or fail with a more sensible
+   error message.  */
+void
+f1 (void)
+{
+  int i;
+
+#pragma omp parallel
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma omp for
+  for (i = 0; i < 3; i++)
+    {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+      ;
+    }
+
+#pragma omp sections
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma omp single
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma omp task
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma omp master
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma omp critical
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma omp ordered
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+}
+
+/* TODO: Some of these should either be allowed or fail with a more sensible
+   error message.  */
+void
+f2 (void)
+{
+#pragma acc parallel
+  {
+#pragma omp parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma acc parallel
+  {
+    int i;
+#pragma omp for		/* { dg-error "may not be nested" } */
+    for (i = 0; i < 3; i++)
+      ;
+  }
+
+#pragma acc parallel
+  {
+#pragma omp sections	/* { dg-error "may not be nested" } */
+    {
+      ;
+    }
+  }
+
+#pragma acc parallel
+  {
+#pragma omp single	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma acc parallel
+  {
+#pragma omp task	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma acc parallel
+  {
+#pragma omp master	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma acc parallel
+  {
+#pragma omp critical	/* { dg-error "may not be nested" } */
+    ;
+  }
+
+#pragma acc parallel
+  {
+    int i;
+#pragma omp atomic write
+    i = 0;		/* { dg-error "may not be nested" } */
+  }
+
+#pragma acc parallel
+  {
+#pragma omp ordered	/* { dg-error "may not be nested" } */
+    ;
+  }
+}
diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
new file mode 100644
index 0000000..6501397
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -0,0 +1,11 @@
+/* TODO: While the OpenACC specification does allow for certain kinds of
+   nesting, we don't support that yet.  */
+void
+f1 (void)
+{
+#pragma acc parallel
+  {
+#pragma acc parallel	/* { dg-error "may not be nested" } */
+    ;
+  }
+}
diff --git gcc/testsuite/c-c++-common/goacc/parallel-1.c gcc/testsuite/c-c++-common/goacc/parallel-1.c
new file mode 100644
index 0000000..cd19527
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/parallel-1.c
@@ -0,0 +1,6 @@
+void
+foo (void)
+{
+#pragma acc parallel
+  foo ();
+}
diff --git gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
new file mode 100644
index 0000000..efc6f14
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
@@ -0,0 +1,6 @@
+void
+foo (void)
+{
+#pragma acc parallel foo	/* { dg-error "expected clause before 'foo'" } */
+  foo ();
+}
diff --git gcc/tree-inline.c gcc/tree-inline.c
index 74f333b..eeb4992 100644
--- gcc/tree-inline.c
+++ gcc/tree-inline.c
@@ -1299,6 +1299,9 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
 	  copy = gimple_build_wce (s1);
 	  break;
 
+	case GIMPLE_OACC_PARALLEL:
+          abort ();
+
 	case GIMPLE_OMP_PARALLEL:
 	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
 	  copy = gimple_build_omp_parallel
@@ -3849,6 +3852,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_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
     case GIMPLE_OMP_CRITICAL:
diff --git gcc/tree-nested.c gcc/tree-nested.c
index dc63ef6..8aba4f4 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1238,6 +1238,9 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	}
       break;
 
+    case GIMPLE_OACC_PARALLEL:
+      abort ();
+
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_suppress = info->suppress_expansion;
@@ -1679,6 +1682,9 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 
   switch (gimple_code (stmt))
     {
+    case GIMPLE_OACC_PARALLEL:
+      abort ();
+
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_suppress = info->suppress_expansion;
@@ -2008,6 +2014,9 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	break;
       }
 
+    case GIMPLE_OACC_PARALLEL:
+      abort ();
+
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       {
@@ -2068,6 +2077,9 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	}
       break;
 
+    case GIMPLE_OACC_PARALLEL:
+      abort ();
+
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_static_chain_added = info->static_chain_added;
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index fe75633..153d01f 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -2346,6 +2346,11 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
       pp_string (buffer, " > ");
       break;
 
+    case OACC_PARALLEL:
+      pp_string (buffer, "#pragma acc parallel");
+      dump_omp_clauses (buffer, OACC_PARALLEL_CLAUSES (node), spc, flags);
+      goto dump_omp_body;
+
     case OMP_PARALLEL:
       pp_string (buffer, "#pragma omp parallel");
       dump_omp_clauses (buffer, OMP_PARALLEL_CLAUSES (node), spc, flags);
diff --git gcc/tree.def gcc/tree.def
index 399b5af..87fec57 100644
--- gcc/tree.def
+++ gcc/tree.def
@@ -1000,8 +1000,15 @@ DEFTREECODE (TARGET_MEM_REF, "target_mem_ref", tcc_reference, 5)
    chain of component references offsetting p by c.  */
 DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2)
 
-/* The ordering of the codes between OMP_PARALLEL and OMP_CRITICAL is
-   exposed to TREE_RANGE_CHECK.  */
+/* OpenACC and OpenMP.  As it is exposed in TREE_RANGE_CHECK invocations, do
+   not change the ordering of these codes.  */
+
+/* OpenACC - #pragma acc parallel [clause1 ... clauseN]
+   Operand 0: OACC_PARALLEL_BODY: Code to be executed in parallel.
+   Operand 1: OACC_PARALLEL_CLAUSES: List of clauses.  */
+
+DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2)
+
 /* OpenMP - #pragma omp parallel [clause1 ... clauseN]
    Operand 0: OMP_PARALLEL_BODY: Code to be executed by all threads.
    Operand 1: OMP_PARALLEL_CLAUSES: List of clauses.  */
diff --git gcc/tree.h gcc/tree.h
index 22a576f..06d94cf 100644
--- gcc/tree.h
+++ gcc/tree.h
@@ -1171,9 +1171,14 @@ extern void protected_set_expr_location (tree, location_t);
 /* OpenMP directive and clause accessors.  */
 
 #define OMP_BODY(NODE) \
-  TREE_OPERAND (TREE_RANGE_CHECK (NODE, OMP_PARALLEL, OMP_CRITICAL), 0)
+  TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_CRITICAL), 0)
 #define OMP_CLAUSES(NODE) \
-  TREE_OPERAND (TREE_RANGE_CHECK (NODE, OMP_PARALLEL, OMP_SINGLE), 1)
+  TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1)
+
+#define OACC_PARALLEL_BODY(NODE) \
+  TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 0)
+#define OACC_PARALLEL_CLAUSES(NODE) \
+  TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1)
 
 #define OMP_PARALLEL_BODY(NODE)    TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 0)
 #define OMP_PARALLEL_CLAUSES(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 1)
diff --git libgomp/Makefile.am libgomp/Makefile.am
index 0b5c097..37b36bd 100644
--- libgomp/Makefile.am
+++ libgomp/Makefile.am
@@ -60,7 +60,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_LDFLAGS)
 libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
 	iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
 	task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
-	time.c fortran.c affinity.c target.c
+	time.c fortran.c affinity.c target.c oacc-parallel.c
 
 nodist_noinst_HEADERS = libgomp_f.h
 nodist_libsubinclude_HEADERS = omp.h openacc.h
diff --git libgomp/Makefile.in libgomp/Makefile.in
index 9ee1bec..bc60253d 100644
--- libgomp/Makefile.in
+++ libgomp/Makefile.in
@@ -96,7 +96,7 @@ am_libgomp_la_OBJECTS = alloc.lo barrier.lo critical.lo env.lo \
 	error.lo iter.lo iter_ull.lo loop.lo loop_ull.lo ordered.lo \
 	parallel.lo sections.lo single.lo task.lo team.lo work.lo \
 	lock.lo mutex.lo proc.lo sem.lo bar.lo ptrlock.lo time.lo \
-	fortran.lo affinity.lo target.lo
+	fortran.lo affinity.lo target.lo oacc-parallel.lo
 libgomp_la_OBJECTS = $(am_libgomp_la_OBJECTS)
 DEFAULT_INCLUDES = -I.@am__isrc@
 depcomp = $(SHELL) $(top_srcdir)/../depcomp
@@ -317,7 +317,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_LDFLAGS)
 libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
 	iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
 	task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
-	time.c fortran.c affinity.c target.c
+	time.c fortran.c affinity.c target.c oacc-parallel.c
 
 nodist_noinst_HEADERS = libgomp_f.h
 nodist_libsubinclude_HEADERS = omp.h openacc.h
@@ -469,6 +469,7 @@ distclean-compile:
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop_ull.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/mutex.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-parallel.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ordered.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/parallel.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/proc.Plo@am__quote@
diff --git libgomp/libgomp.map libgomp/libgomp.map
index f094ed2..2b64d05 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -232,4 +232,6 @@ OACC_2.0 {
 };
 
 GOACC_2.0 {
+  global:
+	GOACC_parallel;
 };
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index 577956a..394f3a8 100644
--- libgomp/libgomp_g.h
+++ libgomp/libgomp_g.h
@@ -214,4 +214,9 @@ extern void GOMP_target_update (int, const void *,
 				size_t, void **, size_t *, unsigned char *);
 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 char *);
+
 #endif /* LIBGOMP_G_H */
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
new file mode 100644
index 0000000..730b83b
--- /dev/null
+++ libgomp/oacc-parallel.c
@@ -0,0 +1,36 @@
+/* Copyright (C) 2013 Free Software Foundation, Inc.
+
+   Contributed by Thomas Schwinge <thomas@codesourcery.com>.
+
+   This file is part of the GNU OpenMP Library (libgomp).
+
+   Libgomp 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.
+
+   Libgomp 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.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This file handles the OpenACC parallel construct.  */
+
+#include "libgomp_g.h"
+
+void
+GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target,
+		size_t mapnum, void **hostaddrs, size_t *sizes,
+		unsigned char *kinds)
+{
+  GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds);
+}
diff --git libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c
new file mode 100644
index 0000000..b9bdffa
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c
@@ -0,0 +1,25 @@
+/* { dg-do run } */
+
+#include "libgomp_g.h"
+
+extern void abort ();
+
+volatile int i;
+
+void
+f (void *data)
+{
+  if (i != -1)
+    abort ();
+  i = 42;
+}
+
+int main(void)
+{
+  i = -1;
+  GOACC_parallel (0, f, (const void *) 0, 0, (void *) 0, (void *) 0, (void *) 0);
+  if (i != 42)
+    abort ();
+
+  return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c/parallel-1.c libgomp/testsuite/libgomp.oacc-c/parallel-1.c
new file mode 100644
index 0000000..b40545d
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+
+extern void abort ();
+
+volatile int i;
+
+int main(void)
+{
+  volatile int j;
+
+  i = -0x42;
+  j = -42;
+#pragma acc parallel
+  {
+    if (i != -0x42 || j != -42)
+      abort ();
+    i = 42;
+    j = 0x42;
+    if (i != 42 || j != 0x42)
+      abort ();
+  }
+  if (i != 42 || j != 0x42)
+    abort ();
+
+  return 0;
+}
