diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index bd46f2e..824ec94 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,20 @@
 2014-02-21  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA.
+	(is_gimple_omp_oacc_specifically): Handle it.
+	* gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
+	* gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise.
+	* omp-low.c (scan_sharing_clauses, scan_omp_target)
+	(expand_omp_target, lower_omp_target, lower_omp_1): Likewise.
+	* gimple.def (GIMPLE_OMP_TARGET): Update comment.
+	* gimple.c (gimple_build_omp_target): Likewise.
+	(gimple_copy): Catch unimplemented case.
+	* tree-inline.c (remap_gimple_stmt): Likewise.
+	* tree-nested.c (convert_nonlocal_reference_stmt)
+	(convert_local_reference_stmt, convert_gimple_call): Likewise.
+	* oacc-builtins.def (BUILT_IN_GOACC_DATA_START)
+	(BUILT_IN_GOACC_DATA_END): New builtins.
+
 	* omp-low.c (scan_sharing_clauses): Catch unexpected occurrences
 	of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP.
 
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index 91a3eb2..ad9369c 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1289,6 +1289,9 @@ 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_DATA:
+      kind = " oacc_data";
+      break;
     default:
       gcc_unreachable ();
     }
diff --git gcc/gimple.c gcc/gimple.c
index 2a967aa..30561b1 100644
--- gcc/gimple.c
+++ gcc/gimple.c
@@ -1051,7 +1051,8 @@ 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.
-   CLAUSES are any of the OMP target construct's clauses.  */
+   KIND is the kind of target region.
+   CLAUSES are any of the construct's clauses.  */
 
 gimple
 gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
@@ -1747,6 +1748,7 @@ gimple_copy (gimple stmt)
 	case GIMPLE_OMP_TASKGROUP:
 	case GIMPLE_OMP_ORDERED:
 	copy_omp_body:
+	  gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
 	  new_seq = gimple_seq_copy (gimple_omp_body (stmt));
 	  gimple_omp_set_body (copy, new_seq);
 	  break;
diff --git gcc/gimple.def gcc/gimple.def
index 2b78c06..ce800bd 100644
--- gcc/gimple.def
+++ gcc/gimple.def
@@ -360,6 +360,7 @@ 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 omp target {,data,update}
    BODY is the sequence of statements inside the target construct
    (NULL for target update).
diff --git gcc/gimple.h gcc/gimple.h
index 0d250ef..b4ee9fa 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -102,6 +102,7 @@ enum gf_mask {
     GF_OMP_TARGET_KIND_REGION	= 0 << 0,
     GF_OMP_TARGET_KIND_DATA	= 1 << 0,
     GF_OMP_TARGET_KIND_UPDATE	= 2 << 0,
+    GF_OMP_TARGET_KIND_OACC_DATA = 3 << 0,
 
     /* True on an GIMPLE_OMP_RETURN statement if the return does not require
        a thread synchronization via some sort of barrier.  The exact barrier
@@ -5684,6 +5685,14 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
     {
     case GIMPLE_OACC_PARALLEL:
       return true;
+    case GIMPLE_OMP_TARGET:
+      switch (gimple_omp_target_kind (stmt))
+	{
+	case GF_OMP_TARGET_KIND_OACC_DATA:
+	  return true;
+	default:
+	  return false;
+	}
     default:
       return false;
     }
diff --git gcc/gimplify.c gcc/gimplify.c
index 9aa9301c..fd4305c 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7023,9 +7023,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
   return GS_ALL_DONE;
 }
 
-/* Gimplify the gross structure of other OpenMP constructs.
-   In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA
-   and OMP_TEAMS.  */
+/* Gimplify the gross structure of several OpenACC or OpenMP constructs.  */
 
 static void
 gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
@@ -7033,12 +7031,17 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
   tree expr = *expr_p;
   gimple stmt;
   gimple_seq body = NULL;
-  enum omp_region_type ort = ORT_WORKSHARE;
+  enum omp_region_type ort;
 
   switch (TREE_CODE (expr))
     {
+    case OACC_DATA:
+      ort = (enum omp_region_type) (ORT_TARGET
+				    | ORT_TARGET_MAP_FORCE);
+      break;
     case OMP_SECTIONS:
     case OMP_SINGLE:
+      ort = ORT_WORKSHARE;
       break;
     case OMP_TARGET:
       ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD);
@@ -7063,9 +7066,21 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 	pop_gimplify_context (NULL);
       if (!(ort & ORT_TARGET_OFFLOAD))
 	{
-	  gimple_seq cleanup = NULL;
-	  tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
+	  enum built_in_function end_ix;
+	  switch (TREE_CODE (expr))
+	    {
+	    case OACC_DATA:
+	      end_ix = BUILT_IN_GOACC_DATA_END;
+	      break;
+	    case OMP_TARGET_DATA:
+	      end_ix = BUILT_IN_GOMP_TARGET_END_DATA;
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	  tree fn = builtin_decl_explicit (end_ix);
 	  g = gimple_build_call (fn, 0);
+	  gimple_seq cleanup = NULL;
 	  gimple_seq_add_stmt (&cleanup, g);
 	  g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
 	  body = NULL;
@@ -7078,6 +7093,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 
   switch (TREE_CODE (expr))
     {
+    case OACC_DATA:
+      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
+				      OACC_DATA_CLAUSES (expr));
+      break;
     case OMP_SECTIONS:
       stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
       break;
@@ -8047,7 +8066,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  break;
 
 	case OACC_KERNELS:
-	case OACC_DATA:
 	case OACC_HOST_DATA:
 	case OACC_DECLARE:
 	case OACC_UPDATE:
@@ -8076,6 +8094,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = gimplify_omp_for (expr_p, pre_p);
 	  break;
 
+	case OACC_DATA:
 	case OMP_SECTIONS:
 	case OMP_SINGLE:
 	case OMP_TARGET:
diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def
index a75e42d..eaf3228 100644
--- gcc/oacc-builtins.def
+++ gcc/oacc-builtins.def
@@ -1,7 +1,7 @@
 /* This file contains the definitions and documentation for the
    OpenACC builtins used in the GNU compiler.
 
-   Copyright (C) 2013 Free Software Foundation, Inc.
+   Copyright (C) 2013-2014 Free Software Foundation, Inc.
 
    Contributed by Thomas Schwinge <thomas@codesourcery.com>.
 
@@ -29,3 +29,7 @@ along with GCC; see the file COPYING3.  If not see
 
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
 		   BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
+		   BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
+		   BT_FN_VOID, ATTR_NOTHROW_LIST)
diff --git gcc/omp-low.c gcc/omp-low.c
index bca4599..6dec687 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1499,6 +1499,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 {
   tree c, decl;
   bool scan_array_reductions = false;
+  bool offloaded;
+  switch (gimple_code (ctx->stmt))
+    {
+    case GIMPLE_OACC_PARALLEL:
+      offloaded = true;
+      break;
+    case GIMPLE_OMP_TARGET:
+      switch (gimple_omp_target_kind (ctx->stmt))
+	{
+	case GF_OMP_TARGET_KIND_REGION:
+	  offloaded = true;
+	  break;
+	case GF_OMP_TARGET_KIND_DATA:
+	case GF_OMP_TARGET_KIND_UPDATE:
+	case GF_OMP_TARGET_KIND_OACC_DATA:
+	  offloaded = false;
+	  break;
+	default:
+	  gcc_unreachable ();
+	}
+      break;
+    default:
+      offloaded = false;
+    }
 
   for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
     {
@@ -1669,11 +1693,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
 	    {
 	      /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
-		 #pragma omp target data, there is nothing to map for
+		 target regions that are not offloaded; there is nothing to map for
 		 those.  */
-	      if (!gimple_code_is_oacc (ctx->stmt)
-		  && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
-		  && !POINTER_TYPE_P (TREE_TYPE (decl)))
+	      if (!offloaded && !POINTER_TYPE_P (TREE_TYPE (decl)))
 		break;
 	    }
 	  if (DECL_P (decl))
@@ -1698,9 +1720,7 @@ 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_code (ctx->stmt) == GIMPLE_OACC_PARALLEL
-		      || (gimple_omp_target_kind (ctx->stmt)
-			  == GF_OMP_TARGET_KIND_REGION))
+		  if (offloaded)
 		    install_var_local (decl, ctx);
 		}
 	    }
@@ -1824,8 +1844,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
 		      || (gimple_omp_target_kind (ctx->stmt)
 			  != GF_OMP_TARGET_KIND_UPDATE));
-	  if (!gimple_code_is_oacc (ctx->stmt)
-	      && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+	  if (!offloaded)
 	    break;
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (DECL_P (decl)
@@ -2340,7 +2359,7 @@ scan_omp_single (gimple stmt, omp_context *outer_ctx)
     layout_type (ctx->record_type);
 }
 
-/* Scan an OpenMP target{, data, update} directive.  */
+/* Scan a GIMPLE_OMP_TARGET.  */
 
 static void
 scan_omp_target (gimple stmt, omp_context *outer_ctx)
@@ -2349,6 +2368,12 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
   tree name;
   int kind = gimple_omp_target_kind (stmt);
 
+  if (kind == GF_OMP_TARGET_KIND_OACC_DATA)
+    {
+      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;
@@ -8218,7 +8243,7 @@ expand_omp_atomic (struct omp_region *region)
 }
 
 
-/* Expand the OpenMP target{, data, update} directive starting at REGION.  */
+/* Expand the GIMPLE_OMP_TARGET starting at REGION.  */
 
 static void
 expand_omp_target (struct omp_region *region)
@@ -8401,12 +8426,23 @@ expand_omp_target (struct omp_region *region)
 
   clauses = gimple_omp_target_clauses (entry_stmt);
 
-  if (kind == GF_OMP_TARGET_KIND_REGION)
-    start_ix = BUILT_IN_GOMP_TARGET;
-  else if (kind == GF_OMP_TARGET_KIND_DATA)
-    start_ix = BUILT_IN_GOMP_TARGET_DATA;
-  else
-    start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+  switch (kind)
+    {
+    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;
+    default:
+      gcc_unreachable ();
+    }
 
   /* By default, the value of DEVICE is -1 (let runtime library choose)
      and there is no conditional.  */
@@ -8414,10 +8450,12 @@ expand_omp_target (struct omp_region *region)
   device = build_int_cst (integer_type_node, -1);
 
   c = find_omp_clause (clauses, OMP_CLAUSE_IF);
+  gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
   if (c)
     cond = OMP_CLAUSE_IF_EXPR (c);
 
   c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
+  gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
   if (c)
     {
       device = OMP_CLAUSE_DEVICE_ID (c);
@@ -8433,6 +8471,7 @@ expand_omp_target (struct omp_region *region)
      (cond ? device : -2).  */
   if (cond)
     {
+      gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
       cond = gimple_boolify (cond);
 
       basic_block cond_bb, then_bb, else_bb;
@@ -8523,7 +8562,9 @@ expand_omp_target (struct omp_region *region)
       gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
       gsi_remove (&gsi, true);
     }
-  if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
+  if ((kind == GF_OMP_TARGET_KIND_DATA
+       || kind == GF_OMP_TARGET_KIND_OACC_DATA)
+      && region->exit)
     {
       gsi = gsi_last_bb (region->exit);
       g = gsi_stmt (gsi);
@@ -10277,7 +10318,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     }
 }
 
-/* Lower the OpenMP target directive in the current statement
+/* Lower the GIMPLE_OMP_TARGET in the current statement
    in GSI_P.  CTX holds context information for the directive.  */
 
 static void
@@ -10298,7 +10339,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       tgt_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
       tgt_body = gimple_bind_body (tgt_bind);
     }
-  else if (kind == GF_OMP_TARGET_KIND_DATA)
+  else if (kind == GF_OMP_TARGET_KIND_DATA
+	   || kind == GF_OMP_TARGET_KIND_OACC_DATA)
     tgt_body = gimple_omp_body (stmt);
   child_fn = ctx->cb.dst_fn;
 
@@ -10322,6 +10364,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case OMP_CLAUSE_MAP_TOFROM:
 	  case OMP_CLAUSE_MAP_POINTER:
 	    break;
+	  case OMP_CLAUSE_MAP_FORCE_ALLOC:
+	  case OMP_CLAUSE_MAP_FORCE_TO:
+	  case OMP_CLAUSE_MAP_FORCE_FROM:
+	  case OMP_CLAUSE_MAP_FORCE_TOFROM:
+	  case OMP_CLAUSE_MAP_FORCE_PRESENT:
+	  case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+	  case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+	    gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA);
+	    break;
 	  default:
 	    gcc_unreachable ();
 	  }
@@ -10330,6 +10381,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
       case OMP_CLAUSE_TO:
       case OMP_CLAUSE_FROM:
+	if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+	  gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
 	var = OMP_CLAUSE_DECL (c);
 	if (!DECL_P (var))
 	  {
@@ -10373,7 +10426,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       lower_omp (&tgt_body, ctx);
       target_nesting_level--;
     }
-  else if (kind == GF_OMP_TARGET_KIND_DATA)
+  else if (kind == GF_OMP_TARGET_KIND_DATA
+	   || kind == GF_OMP_TARGET_KIND_OACC_DATA)
     lower_omp (&tgt_body, ctx);
 
   if (kind == GF_OMP_TARGET_KIND_REGION)
@@ -10400,9 +10454,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       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 tkind_type;
+      int talign_shift;
+      switch (kind)
+	{
+	case GF_OMP_TARGET_KIND_REGION:
+	case GF_OMP_TARGET_KIND_DATA:
+	case GF_OMP_TARGET_KIND_UPDATE:
+	  tkind_type = unsigned_char_type_node;
+	  talign_shift = 3;
+	  break;
+	case GF_OMP_TARGET_KIND_OACC_DATA:
+	  tkind_type = short_unsigned_type_node;
+	  talign_shift = 8;
+	  break;
+	default:
+	  gcc_unreachable ();
+	}
       TREE_VEC_ELT (t, 2)
-	= create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
-						  map_cnt),
+	= create_tmp_var (build_array_type_nelts (tkind_type, map_cnt),
 			  ".omp_data_kinds");
       DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
       TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
@@ -10515,7 +10585,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    if (TREE_CODE (s) != INTEGER_CST)
 	      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
 
-	    unsigned char tkind = 0;
+	    unsigned HOST_WIDE_INT tkind;
 	    switch (OMP_CLAUSE_CODE (c))
 	      {
 	      case OMP_CLAUSE_MAP:
@@ -10530,14 +10600,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	      default:
 		gcc_unreachable ();
 	      }
-	    unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+	    gcc_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift));
+	    unsigned HOST_WIDE_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;
+	    tkind |= talign << talign_shift;
+	    gcc_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
-				    build_int_cst (unsigned_char_type_node,
-						   tkind));
+				    build_int_cstu (tkind_type, tkind));
 	    if (nc && nc != c)
 	      c = nc;
 	  }
@@ -10589,7 +10660,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       gimple_seq_add_seq (&new_body, tgt_body);
       new_body = maybe_catch_exception (new_body);
     }
-  else if (kind == GF_OMP_TARGET_KIND_DATA)
+  else if (kind == GF_OMP_TARGET_KIND_DATA
+	   || kind == GF_OMP_TARGET_KIND_OACC_DATA)
     new_body = tgt_body;
   if (kind != GF_OMP_TARGET_KIND_UPDATE)
     {
@@ -10810,6 +10882,8 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     case GIMPLE_OMP_TARGET:
       ctx = maybe_lookup_ctx (stmt);
       gcc_assert (ctx);
+      if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA)
+	gcc_assert (!ctx->cancellable);
       lower_omp_target (gsi_p, ctx);
       break;
     case GIMPLE_OMP_TEAMS:
diff --git gcc/tree-inline.c gcc/tree-inline.c
index 99903333..61c1cc8 100644
--- gcc/tree-inline.c
+++ gcc/tree-inline.c
@@ -1397,6 +1397,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
 	  break;
 
 	case GIMPLE_OMP_TARGET:
+	  gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
 	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
 	  copy = gimple_build_omp_target
 		   (s1, gimple_omp_target_kind (stmt),
diff --git gcc/tree-nested.c gcc/tree-nested.c
index 8933d02..afa7abb 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1307,6 +1307,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_TARGET:
+      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       save_suppress = info->suppress_expansion;
       convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
       walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op,
@@ -1769,6 +1770,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_TARGET:
+      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       save_suppress = info->suppress_expansion;
       convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
       walk_body (convert_local_reference_stmt, convert_local_reference_op,
@@ -2184,6 +2186,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
     case GIMPLE_OMP_TASKGROUP:
     case GIMPLE_OMP_ORDERED:
     case GIMPLE_OMP_CRITICAL:
+      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
       break;
 
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 3dffde4..5c15656 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2014-02-21  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start.
+	* libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes.
+	* oacc-parallel.c (GOACC_data_start, GOACC_data_end): New
+	functions.
+
 2014-02-20  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* target.c (gomp_load_plugin_for_device): Don't call dlcose if
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 2b64d05..cb52e45 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -233,5 +233,7 @@ OACC_2.0 {
 
 GOACC_2.0 {
   global:
+	GOACC_data_end;
+	GOACC_data_start;
 	GOACC_parallel;
 };
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index 7c24317..b9083a5 100644
--- libgomp/libgomp_g.h
+++ libgomp/libgomp_g.h
@@ -218,5 +218,8 @@ extern void GOMP_teams (unsigned int, unsigned int);
 
 extern void GOACC_parallel (int, void (*) (void *), const void *,
 			    size_t, void **, size_t *, unsigned short *);
+extern void GOACC_data_start (int, const void *,
+			      size_t, void **, size_t *, unsigned short *);
+extern void GOACC_data_end (void);
 
 #endif /* LIBGOMP_G_H */
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index bf7b74c..3ac7e39 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -1,4 +1,4 @@
-/* Copyright (C) 2013 Free Software Foundation, Inc.
+/* Copyright (C) 2013-2014 Free Software Foundation, Inc.
 
    Contributed by Thomas Schwinge <thomas@codesourcery.com>.
 
@@ -23,7 +23,7 @@
    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
    <http://www.gnu.org/licenses/>.  */
 
-/* This file handles the OpenACC parallel construct.  */
+/* This file handles the OpenACC data and parallel constructs.  */
 
 #include "libgomp.h"
 #include "libgomp_g.h"
@@ -51,3 +51,33 @@ GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target,
     }
   GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds_);
 }
+
+
+void
+GOACC_data_start (int device, const void *openmp_target, size_t mapnum,
+		  void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+  unsigned char kinds_[mapnum];
+  size_t i;
+
+  /* TODO.  Eventually, we'll be interpreting all mapping kinds according to
+     the OpenACC semantics; for now we're re-using what is implemented for
+     OpenMP.  */
+  for (i = 0; i < mapnum; ++i)
+    {
+      unsigned char kind = kinds[i];
+      unsigned char align = kinds[i] >> 8;
+      if (kind > 4)
+	gomp_fatal ("memory mapping kind %x for %zd is not yet supported",
+		    kind, i);
+
+      kinds_[i] = kind | align << 3;
+    }
+  GOMP_target_data (device, openmp_target, mapnum, hostaddrs, sizes, kinds_);
+}
+
+void
+GOACC_data_end (void)
+{
+  GOMP_target_end_data ();
+}
